diff -urpN magnum-jumbo/src/opencl/phpass_kernel.cl magnum-jumbo_phpass-ocl//src/opencl/phpass_kernel.cl
--- magnum-jumbo/src/opencl/phpass_kernel.cl	2012-08-21 17:35:10.087759411 +0000
+++ magnum-jumbo_phpass-ocl//src/opencl/phpass_kernel.cl	2012-08-21 17:39:41.000000000 +0000
@@ -1,196 +1,295 @@
 /*
-* This software is Copyright (c) 2011,2012 Lukas Odzioba <ukasz at openwall dot net>
+* This software is Copyright (c) 2011-2012 Lukas Odzioba <ukasz at openwall dot net> 
 * and it is hereby released to the general public under the following terms:
 * Redistribution and use in source and binary forms, with or without modification, are permitted.
 */
-#define PLAINTEXT_LENGTH	15
+#define PLAINTEXT_LENGTH        15
+#define SALT_SIZE               8
+//#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : disable
+#pragma OPENCL EXTENSION cl_amd_media_ops : enable
+
 typedef struct {
-	unsigned char v[PLAINTEXT_LENGTH];
-	unsigned char length;
+	uchar v[PLAINTEXT_LENGTH];
+	uchar length;
 } phpass_password;
 
 typedef struct {
-	unsigned int v[4];
-} phpass_hash;
+	uchar salt[8];
+	uint hash[4];
+	uint rounds;
+} phpass_salt;
+
+typedef struct {
+	uchar cracked;
+} phpass_crack;
+
 
-#define ROTATE_LEFT(x, s)	rotate(x,(unsigned int)s)
+//#define H(x, y, z)              ((x) ^ (y) ^ (z))
+//#define I(x, y, z)              ((y) ^ ((x) | (~z)))
+
+//#define ROTATE_LEFT(x, s)       rotate(x,(uint)s)
+#define ROTATE_LEFT(x, s)  amd_bitalign(x, x, (uint)(32 - s))
+//#define ROTATE_LEFT(a,s) (a<<s)+(a>>(32-s))
+#define F(x, y, z) bitselect((z), (y), (x))
+#define G(x, y, z) bitselect((y), (x), (z))
+
+
+//#define F(x, y, z)                    ((z) ^ ((x) & ((y) ^ (z))))
+//#define G(x, y, z)                    ((y) ^ ((z) & ((x) ^ (y))))
+#define H(x, y, z)			((x) ^ (y) ^ (z))
+#define I(x, y, z)			((y) ^ ((x) | ~(z)))
 
-#define F(x, y, z)		bitselect((z), (y), (x))
-#define G(x, y, z)		bitselect((y), (x), (z))
-#define H(x, y, z)		((x) ^ (y) ^ (z))
-#define I(x, y, z)		((y) ^ ((x) | (~z)))
 
 #define FF(a, b, c, d, x, s, ac) \
-  {(a) += F ((b), (c), (d)) + (x) + (uint32_t)(ac); \
-   (a) = ROTATE_LEFT ((a), (s)); \
-   (a) += (b); \
-  }
+   (a) = ROTATE_LEFT ( a + x + ac+F ((b), (c), (d)), (s)) + (b);
+
 #define GG(a, b, c, d, x, s, ac) \
-  {(a) += G ((b), (c), (d)) + (x) + (uint32_t)(ac); \
-   (a) = ROTATE_LEFT ((a), (s)); \
-   (a) += (b); \
-  }
+   (a) = ROTATE_LEFT ( a + x + ac +G ((b), (c), (d)), (s)) + (b);
+
 #define HH(a, b, c, d, x, s, ac) \
-  {(a) += H ((b), (c), (d)) + (x) + (uint32_t)(ac); \
-   (a) = ROTATE_LEFT ((a), (s)); \
-   (a) += (b); \
-  }
+   (a) = ROTATE_LEFT ( a + x + ac +H ((b), (c), (d)) , (s)) + (b);
+
 #define II(a, b, c, d, x, s, ac) \
-  {(a) += I ((b), (c), (d)) + (x) + (uint32_t)(ac); \
-   (a) = ROTATE_LEFT ((a), (s)); \
-   (a) += (b); \
-  }
-
-#define S11 7
-#define S12 12
-#define S13 17
-#define S14 22
-#define S21 5
-#define S22 9
-#define S23 14
-#define S24 20
-#define S31 4
-#define S32 11
-#define S33 16
-#define S34 23
-#define S41 6
-#define S42 10
-#define S43 15
-#define S44 21
-#define uint32_t unsigned int
-#define SALT_SIZE		8
-
-#define AC1				0xd76aa477
-#define AC2pCd				0xf8fa0bcc
-#define AC3pCc				0xbcdb4dd9
-#define AC4pCb				0xb18b7a77
-#define MASK1				0x77777777
+   (a) = ROTATE_LEFT ( a + x + ac +I ((b), (c), (d)) , (s)) + (b);
 
-inline void md5(char len,__private uint32_t * internal_ret,__private uint32_t * x)
-{
-	x[len / 4] |= (((uint32_t) 0x80) << ((len & 0x3) << 3));
-	uint32_t x14 = len << 3;
+#define S11                             7
+#define S12                             12
+#define S13                             17
+#define S14                             22
+#define S21                             5
+#define S22                             9
+#define S23                             14
+#define S24                             20
+#define S31                             4
+#define S32                             11
+#define S33                             16
+#define S34                             23
+#define S41                             6
+#define S42                             10
+#define S43                             15
+#define S44                             21
+
+#define AC1                             (uint8) 0xd76aa477
+#define AC2pCd                          (uint8) 0xf8fa0bcc
+#define AC3pCc                          (uint8) 0xbcdb4dd9
+#define AC4pCb                          (uint8) 0xb18b7a77
+#define MASK1                           (uint8) 0x77777777
+
+#define FF2(v, w, x, y, s, ac) \
+ v = ROTATE_LEFT(v + ac + F(w, x, y), s) + w;
+
+#define GG2(v, w, x, y, s, ac)  \
+ v = ROTATE_LEFT(v + ac + G(w, x, y), s) + w;
+
+#define HH2(v, w, x, y, s, ac)  \
+ v = ROTATE_LEFT(v + ac + H(w, x, y), s) + w;
+
+#define II2(v, w, x, y, s, ac)  \
+ v = ROTATE_LEFT(v + ac + I(w, x, y), s) + w;
 
-	uint32_t a = 0x67452301;
-	uint32_t b = 0xefcdab89;
-	uint32_t c = 0x98badcfe;
-	uint32_t d = 0x10325476;
-
-	FF(a, b, c, d, x[0], S11, 0xd76aa478);
-	FF(d, a, b, c, x[1], S12, 0xe8c7b756);
-	FF(c, d, a, b, x[2], S13, 0x242070db);
-	FF(b, c, d, a, x[3], S14, 0xc1bdceee);
-	FF(a, b, c, d, x[4], S11, 0xf57c0faf);
-	FF(d, a, b, c, x[5], S12, 0x4787c62a);
-	FF(c, d, a, b, x[6], S13, 0xa8304613);
-	FF(b, c, d, a, x[7], S14, 0xfd469501);
-	FF(a, b, c, d, 0, S11, 0x698098d8);
-	FF(d, a, b, c, 0, S12, 0x8b44f7af);
-	FF(c, d, a, b, 0, S13, 0xffff5bb1);
-	FF(b, c, d, a, 0, S14, 0x895cd7be);
-	FF(a, b, c, d, 0, S11, 0x6b901122);
-	FF(d, a, b, c, 0, S12, 0xfd987193);
-	FF(c, d, a, b, x14, S13, 0xa679438e);
-	FF(b, c, d, a, 0, S14, 0x49b40821);
-
-	GG(a, b, c, d, x[1], S21, 0xf61e2562);
-	GG(d, a, b, c, x[6], S22, 0xc040b340);
-	GG(c, d, a, b, 0, S23, 0x265e5a51);
-	GG(b, c, d, a, x[0], S24, 0xe9b6c7aa);
-	GG(a, b, c, d, x[5], S21, 0xd62f105d);
-	GG(d, a, b, c, 0, S22, 0x2441453);
-	GG(c, d, a, b, 0, S23, 0xd8a1e681);
-	GG(b, c, d, a, x[4], S24, 0xe7d3fbc8);
-	GG(a, b, c, d, 0, S21, 0x21e1cde6);
-	GG(d, a, b, c, x14, S22, 0xc33707d6);
-	GG(c, d, a, b, x[3], S23, 0xf4d50d87);
-	GG(b, c, d, a, 0, S24, 0x455a14ed);
-	GG(a, b, c, d, 0, S21, 0xa9e3e905);
-	GG(d, a, b, c, x[2], S22, 0xfcefa3f8);
-	GG(c, d, a, b, x[7], S23, 0x676f02d9);
-	GG(b, c, d, a, 0, S24, 0x8d2a4c8a);
-
-	HH(a, b, c, d, x[5], S31, 0xfffa3942);
-	HH(d, a, b, c, 0, S32, 0x8771f681);
-	HH(c, d, a, b, 0, S33, 0x6d9d6122);
-	HH(b, c, d, a, x14, S34, 0xfde5380c);
-	HH(a, b, c, d, x[1], S31, 0xa4beea44);
-	HH(d, a, b, c, x[4], S32, 0x4bdecfa9);
-	HH(c, d, a, b, x[7], S33, 0xf6bb4b60);
-	HH(b, c, d, a, 0, S34, 0xbebfbc70);
-	HH(a, b, c, d, 0, S31, 0x289b7ec6);
-	HH(d, a, b, c, x[0], S32, 0xeaa127fa);
-	HH(c, d, a, b, x[3], S33, 0xd4ef3085);
-	HH(b, c, d, a, x[6], S34, 0x4881d05);
-	HH(a, b, c, d, 0, S31, 0xd9d4d039);
-	HH(d, a, b, c, 0, S32, 0xe6db99e5);
-	HH(c, d, a, b, 0, S33, 0x1fa27cf8);
-	HH(b, c, d, a, x[2], S34, 0xc4ac5665);
-
-	II(a, b, c, d, x[0], S41, 0xf4292244);
-	II(d, a, b, c, x[7], S42, 0x432aff97);
-	II(c, d, a, b, x14, S43, 0xab9423a7);
-	II(b, c, d, a, x[5], S44, 0xfc93a039);
-	II(a, b, c, d, 0, S41, 0x655b59c3);
-	II(d, a, b, c, x[3], S42, 0x8f0ccc92);
-	II(c, d, a, b, 0, S43, 0xffeff47d);
-	II(b, c, d, a, x[1], S44, 0x85845dd1);
-	II(a, b, c, d, 0, S41, 0x6fa87e4f);
-	II(d, a, b, c, 0, S42, 0xfe2ce6e0);
-	II(c, d, a, b, x[6], S43, 0xa3014314);
-	II(b, c, d, a, 0, S44, 0x4e0811a1);
-	II(a, b, c, d, x[4], S41, 0xf7537e82);
-	II(d, a, b, c, 0, S42, 0xbd3af235);
-	II(c, d, a, b, x[2], S43, 0x2ad7d2bb);
-	II(b, c, d, a, 0, S44, 0xeb86d391);
-
-	internal_ret[0] = a + 0x67452301;
-	internal_ret[1] = b + 0xefcdab89;
-	internal_ret[2] = c + 0x98badcfe;
-	internal_ret[3] = d + 0x10325476;
+
+
+inline void clear_ctx(__private uint8 * x)
+{
+	uint8 zero = (uint8) (0, 0, 0, 0, 0, 0, 0, 0);
+	for (int i = 0; i < 8; i++)
+		x[i] = zero;
 }
 
-inline void clear_ctx(__private uint32_t * x)
+inline void clean_ctx(__private uint * x)
 {
-	int i;
-	for (i = 0; i < 8; i++)
-		*x++ = 0;
+	for (int i = 0; i < 8; i++)
+		x[i] = 0;
 }
 
+
 __kernel void phpass
-    (   __global    const   phpass_password*    data
-    ,   __global            phpass_hash*    	data_out
-    ,   __global    const   char* 		setting
-    )
-{
-	uint32_t x[8];
-	clear_ctx(x);
-	uint32_t a, b, c, d, x0, x1, x2, x3, x4, x5, x6, x7;
+    (__global const phpass_password * data, __global phpass_crack * data_out,
+    __global const phpass_salt * salt) {
+	uint8 x[8], length;
+	uint sx[8], i, idx = get_global_id(0);
+	uint count = salt->rounds;
 
-	uint32_t idx = get_global_id(0);
+	clear_ctx(x);
 
-	__global const char *password = (__global const char*) data[idx].v;
-	int length, count, i;
-	__private unsigned char *buff = (unsigned char *) x;
+	__global const uchar *password0 = data[idx * 8 + 0].v;
+	__global const uchar *password1 = data[idx * 8 + 1].v;
+	__global const uchar *password2 = data[idx * 8 + 2].v;
+	__global const uchar *password3 = data[idx * 8 + 3].v;
+	__global const uchar *password4 = data[idx * 8 + 4].v;
+	__global const uchar *password5 = data[idx * 8 + 5].v;
+	__global const uchar *password6 = data[idx * 8 + 6].v;
+	__global const uchar *password7 = data[idx * 8 + 7].v;
+
+
+	length.s0 = (uint) data[idx * 8 + 0].length;
+	length.s1 = (uint) data[idx * 8 + 1].length;
+	length.s2 = (uint) data[idx * 8 + 2].length;
+	length.s3 = (uint) data[idx * 8 + 3].length;
+	length.s4 = (uint) data[idx * 8 + 4].length;
+	length.s5 = (uint) data[idx * 8 + 5].length;
+	length.s6 = (uint) data[idx * 8 + 6].length;
+	length.s7 = (uint) data[idx * 8 + 7].length;
+
+	uint8 a, b, c, d, x0, x1, x2, x3, x4, x5, x6, x7, x14;
 
-	length = data[idx].length;
 
-	for (i = 0; i < 8; i++)
-		buff[i] = setting[i];
 
-	for (i = 8; i < 8 + length; i++) {
-		buff[i] = password[i - 8];
-	}
+	__private uchar *buff = (uchar *) sx;
+
+#define K1(q)\
+                clean_ctx(sx);\
+                for (i = 0; i < 8; i++)\
+                        buff[i] = salt->salt[i];\
+                for (i = 8; i < 8 + length.s##q; i++)\
+                        buff[i] = password##q[i - 8];\
+                for ( i = 0; i < 8; i++)\
+                        x[i].s##q=sx[i];
+	K1(0);
+	K1(1);
+	K1(2);
+	K1(3);
+	K1(4);
+	K1(5);
+	K1(6);
+	K1(7);
+#undef K1
+
+
+	uint8 len = length + (uint8) (8);
+
+
+	x[len.s0 / 4].s0 |= (((uint) 0x80) << ((len.s0 & 0x3) << 3));
+	x[len.s1 / 4].s1 |= (((uint) 0x80) << ((len.s1 & 0x3) << 3));
+	x[len.s2 / 4].s2 |= (((uint) 0x80) << ((len.s2 & 0x3) << 3));
+	x[len.s3 / 4].s3 |= (((uint) 0x80) << ((len.s3 & 0x3) << 3));
+	x[len.s4 / 4].s4 |= (((uint) 0x80) << ((len.s4 & 0x3) << 3));
+	x[len.s5 / 4].s5 |= (((uint) 0x80) << ((len.s5 & 0x3) << 3));
+	x[len.s6 / 4].s6 |= (((uint) 0x80) << ((len.s6 & 0x3) << 3));
+	x[len.s7 / 4].s7 |= (((uint) 0x80) << ((len.s7 & 0x3) << 3));
 
-	md5(8 + length, x, x);
-	count = 1 << setting[SALT_SIZE+3];
-	for (i = 16; i < 16 + length; i++)
-		buff[i] = password[i - 16];
+	x0 = x[0];
+	x1 = x[1];
+	x2 = x[2];
+	x3 = x[3];
+	x4 = x[4];
+	x5 = x[5];
+	x6 = x[6];
+	x7 = x[7];
+	x14 = len << 3;
 
-	uint32_t len = 16 + length;
-	uint32_t x14 = len << 3;
+	b = (uint8) 0xefcdab89;
+	c = (uint8) 0x98badcfe;
+	d = (uint8) 0x10325476;
+
+	a = AC1 + x0;
+	a = ROTATE_LEFT(a, S11);
+	a += b;
+	d = (c ^ (a & MASK1)) + x1 + AC2pCd;
+	d = ROTATE_LEFT(d, S12);
+	d += a;
+	c = F(d, a, b) + x2 + AC3pCc;
+	c = ROTATE_LEFT(c, S13);
+	c += d;
+	b = F(c, d, a) + x3 + AC4pCb;
+	b = ROTATE_LEFT(b, S14);
+	b += c;
+	FF(a, b, c, d, x4, S11, (uint8) 0xf57c0faf);
+	FF(d, a, b, c, x5, S12, (uint8) 0x4787c62a);
+	FF(c, d, a, b, x6, S13, (uint8) 0xa8304613);
+	FF(b, c, d, a, x7, S14, (uint8) 0xfd469501);
+	FF2(a, b, c, d, S11, (uint8) 0x698098d8);
+	FF2(d, a, b, c, S12, (uint8) 0x8b44f7af);
+	FF2(c, d, a, b, S13, (uint8) 0xffff5bb1);
+	FF2(b, c, d, a, S14, (uint8) 0x895cd7be);
+	FF2(a, b, c, d, S11, (uint8) 0x6b901122);
+	FF2(d, a, b, c, S12, (uint8) 0xfd987193);
+	FF(c, d, a, b, x14, S13, (uint8) 0xa679438e);
+	FF2(b, c, d, a, S14, (uint8) 0x49b40821);
+
+	GG(a, b, c, d, x1, S21, (uint8) 0xf61e2562);
+	GG(d, a, b, c, x6, S22, (uint8) 0xc040b340);
+	GG2(c, d, a, b, S23, (uint8) 0x265e5a51);
+	GG(b, c, d, a, x0, S24, (uint8) 0xe9b6c7aa);
+	GG(a, b, c, d, x5, S21, (uint8) 0xd62f105d);
+	GG2(d, a, b, c, S22, (uint8) 0x2441453);
+	GG2(c, d, a, b, S23, (uint8) 0xd8a1e681);
+	GG(b, c, d, a, x4, S24, (uint8) 0xe7d3fbc8);
+	GG2(a, b, c, d, S21, (uint8) 0x21e1cde6);
+	GG(d, a, b, c, x14, S22, (uint8) 0xc33707d6);
+	GG(c, d, a, b, x3, S23, (uint8) 0xf4d50d87);
+	GG2(b, c, d, a, S24, (uint8) 0x455a14ed);
+	GG2(a, b, c, d, S21, (uint8) 0xa9e3e905);
+	GG(d, a, b, c, x2, S22, (uint8) 0xfcefa3f8);
+	GG(c, d, a, b, x7, S23, (uint8) 0x676f02d9);
+	GG2(b, c, d, a, S24, (uint8) 0x8d2a4c8a);
+
+	HH(a, b, c, d, x5, S31, (uint8) 0xfffa3942);
+	HH2(d, a, b, c, S32, (uint8) 0x8771f681);
+	HH2(c, d, a, b, S33, (uint8) 0x6d9d6122);
+	HH(b, c, d, a, x14, S34, (uint8) 0xfde5380c);
+	HH(a, b, c, d, x1, S31, (uint8) 0xa4beea44);
+	HH(d, a, b, c, x4, S32, (uint8) 0x4bdecfa9);
+	HH(c, d, a, b, x7, S33, (uint8) 0xf6bb4b60);
+	HH2(b, c, d, a, S34, (uint8) 0xbebfbc70);
+	HH2(a, b, c, d, S31, (uint8) 0x289b7ec6);
+	HH(d, a, b, c, x0, S32, (uint8) 0xeaa127fa);
+	HH(c, d, a, b, x3, S33, (uint8) 0xd4ef3085);
+	HH(b, c, d, a, x6, S34, (uint8) 0x4881d05);
+	HH2(a, b, c, d, S31, (uint8) 0xd9d4d039);
+	HH2(d, a, b, c, S32, (uint8) 0xe6db99e5);
+	HH2(c, d, a, b, S33, (uint8) 0x1fa27cf8);
+	HH(b, c, d, a, x2, S34, (uint8) 0xc4ac5665);
+
+	II(a, b, c, d, x0, S41, (uint8) 0xf4292244);
+	II(d, a, b, c, x7, S42, (uint8) 0x432aff97);
+	II(c, d, a, b, x14, S43, (uint8) 0xab9423a7);
+	II(b, c, d, a, x5, S44, (uint8) 0xfc93a039);
+	II2(a, b, c, d, S41, (uint8) 0x655b59c3);
+	II(d, a, b, c, x3, S42, (uint8) 0x8f0ccc92);
+	II2(c, d, a, b, S43, (uint8) 0xffeff47d);
+	II(b, c, d, a, x1, S44, (uint8) 0x85845dd1);
+	II2(a, b, c, d, S41, (uint8) 0x6fa87e4f);
+	II2(d, a, b, c, S42, (uint8) 0xfe2ce6e0);
+	II(c, d, a, b, x6, S43, (uint8) 0xa3014314);
+	II2(b, c, d, a, S44, (uint8) 0x4e0811a1);
+	II(a, b, c, d, x4, S41, (uint8) 0xf7537e82);
+	II2(d, a, b, c, S42, (uint8) 0xbd3af235);
+	II(c, d, a, b, x2, S43, (uint8) 0x2ad7d2bb);
+	II2(b, c, d, a, S44, (uint8) 0xeb86d391);
+
+	x[0] = a + (uint8) 0x67452301;
+	x[1] = b + (uint8) 0xefcdab89;
+	x[2] = c + (uint8) 0x98badcfe;
+	x[3] = d + (uint8) 0x10325476;
+
+#define K2(q)\
+                clean_ctx(sx);\
+                for(i=0;i<length.s##q;i++)\
+                        buff[i]=password##q[i];\
+                for(i=0;i<4;i++)\
+                        x[i+4].s##q=sx[i];
+	K2(0);
+	K2(1);
+	K2(2);
+	K2(3);
+	K2(4);
+	K2(5);
+	K2(6);
+	K2(7);
+#undef K2
+
+	len = (uint8) (16) + length;
+	x[len.s0 / 4].s0 |= (((uint) 0x80) << ((len.s0 & 0x3) << 3));
+	x[len.s1 / 4].s1 |= (((uint) 0x80) << ((len.s1 & 0x3) << 3));
+	x[len.s2 / 4].s2 |= (((uint) 0x80) << ((len.s2 & 0x3) << 3));
+	x[len.s3 / 4].s3 |= (((uint) 0x80) << ((len.s3 & 0x3) << 3));
+	x[len.s4 / 4].s4 |= (((uint) 0x80) << ((len.s4 & 0x3) << 3));
+	x[len.s5 / 4].s5 |= (((uint) 0x80) << ((len.s5 & 0x3) << 3));
+	x[len.s6 / 4].s6 |= (((uint) 0x80) << ((len.s6 & 0x3) << 3));
+	x[len.s7 / 4].s7 |= (((uint) 0x80) << ((len.s7 & 0x3) << 3));
 
-	x[len / 4] |= ((0x80) << ((len & 0x3) << 3));
+	x14 = len << 3;
 	x0 = x[0];
 	x1 = x[1];
 	x2 = x[2];
@@ -199,11 +298,11 @@ __kernel void phpass
 	x5 = x[5];
 	x6 = x[6];
 	x7 = x[7];
-do {
+	do {
 
-		b = 0xefcdab89;
-		c = 0x98badcfe;
-		d = 0x10325476;
+		b = (uint8) 0xefcdab89;
+		c = (uint8) 0x98badcfe;
+		d = (uint8) 0x10325476;
 
 		a = AC1 + x0;
 		a = ROTATE_LEFT(a, S11);
@@ -217,79 +316,92 @@ do {
 		b = F(c, d, a) + x3 + AC4pCb;
 		b = ROTATE_LEFT(b, S14);
 		b += c;
-		FF(a, b, c, d, x4, S11, 0xf57c0faf);
-		FF(d, a, b, c, x5, S12, 0x4787c62a);
-		FF(c, d, a, b, x6, S13, 0xa8304613);
-		FF(b, c, d, a, x7, S14, 0xfd469501);
-		FF(a, b, c, d, 0, S11, 0x698098d8);
-		FF(d, a, b, c, 0, S12, 0x8b44f7af);
-		FF(c, d, a, b, 0, S13, 0xffff5bb1);
-		FF(b, c, d, a, 0, S14, 0x895cd7be);
-		FF(a, b, c, d, 0, S11, 0x6b901122);
-		FF(d, a, b, c, 0, S12, 0xfd987193);
-		FF(c, d, a, b, x14, S13, 0xa679438e);
-		FF(b, c, d, a, 0, S14, 0x49b40821);
-
-		GG(a, b, c, d, x1, S21, 0xf61e2562);
-		GG(d, a, b, c, x6, S22, 0xc040b340);
-		GG(c, d, a, b, 0, S23, 0x265e5a51);
-		GG(b, c, d, a, x0, S24, 0xe9b6c7aa);
-		GG(a, b, c, d, x5, S21, 0xd62f105d);
-		GG(d, a, b, c, 0, S22, 0x2441453);
-		GG(c, d, a, b, 0, S23, 0xd8a1e681);
-		GG(b, c, d, a, x4, S24, 0xe7d3fbc8);
-		GG(a, b, c, d, 0, S21, 0x21e1cde6);
-		GG(d, a, b, c, x14, S22, 0xc33707d6);
-		GG(c, d, a, b, x3, S23, 0xf4d50d87);
-		GG(b, c, d, a, 0, S24, 0x455a14ed);
-		GG(a, b, c, d, 0, S21, 0xa9e3e905);
-		GG(d, a, b, c, x2, S22, 0xfcefa3f8);
-		GG(c, d, a, b, x7, S23, 0x676f02d9);
-		GG(b, c, d, a, 0, S24, 0x8d2a4c8a);
-
-		HH(a, b, c, d, x5, S31, 0xfffa3942);
-		HH(d, a, b, c, 0, S32, 0x8771f681);
-		HH(c, d, a, b, 0, S33, 0x6d9d6122);
-		HH(b, c, d, a, x14, S34, 0xfde5380c);
-		HH(a, b, c, d, x1, S31, 0xa4beea44);
-		HH(d, a, b, c, x4, S32, 0x4bdecfa9);
-		HH(c, d, a, b, x7, S33, 0xf6bb4b60);
-		HH(b, c, d, a, 0, S34, 0xbebfbc70);
-		HH(a, b, c, d, 0, S31, 0x289b7ec6);
-		HH(d, a, b, c, x0, S32, 0xeaa127fa);
-		HH(c, d, a, b, x3, S33, 0xd4ef3085);
-		HH(b, c, d, a, x6, S34, 0x4881d05);
-		HH(a, b, c, d, 0, S31, 0xd9d4d039);
-		HH(d, a, b, c, 0, S32, 0xe6db99e5);
-		HH(c, d, a, b, 0, S33, 0x1fa27cf8);
-		HH(b, c, d, a, x2, S34, 0xc4ac5665);
-
-		II(a, b, c, d, x0, S41, 0xf4292244);
-		II(d, a, b, c, x7, S42, 0x432aff97);
-		II(c, d, a, b, x14, S43, 0xab9423a7);
-		II(b, c, d, a, x5, S44, 0xfc93a039);
-		II(a, b, c, d, 0, S41, 0x655b59c3);
-		II(d, a, b, c, x3, S42, 0x8f0ccc92);
-		II(c, d, a, b, 0, S43, 0xffeff47d);
-		II(b, c, d, a, x1, S44, 0x85845dd1);
-		II(a, b, c, d, 0, S41, 0x6fa87e4f);
-		II(d, a, b, c, 0, S42, 0xfe2ce6e0);
-		II(c, d, a, b, x6, S43, 0xa3014314);
-		II(b, c, d, a, 0, S44, 0x4e0811a1);
-		II(a, b, c, d, x4, S41, 0xf7537e82);
-		II(d, a, b, c, 0, S42, 0xbd3af235);
-		II(c, d, a, b, x2, S43, 0x2ad7d2bb);
-		II(b, c, d, a, 0, S44, 0xeb86d391);
-
-		x0 = a + 0x67452301;
-		x1 = b + 0xefcdab89;
-		x2 = c + 0x98badcfe;
-		x3 = d + 0x10325476;
-
+		FF(a, b, c, d, x4, S11, (uint8) 0xf57c0faf);
+		FF(d, a, b, c, x5, S12, (uint8) 0x4787c62a);
+		FF(c, d, a, b, x6, S13, (uint8) 0xa8304613);
+		FF(b, c, d, a, x7, S14, (uint8) 0xfd469501);
+		FF2(a, b, c, d, S11, (uint8) 0x698098d8);
+		FF2(d, a, b, c, S12, (uint8) 0x8b44f7af);
+		FF2(c, d, a, b, S13, (uint8) 0xffff5bb1);
+		FF2(b, c, d, a, S14, (uint8) 0x895cd7be);
+		FF2(a, b, c, d, S11, (uint8) 0x6b901122);
+		FF2(d, a, b, c, S12, (uint8) 0xfd987193);
+		FF(c, d, a, b, x14, S13, (uint8) 0xa679438e);
+		FF2(b, c, d, a, S14, (uint8) 0x49b40821);
+
+		GG(a, b, c, d, x1, S21, (uint8) 0xf61e2562);
+		GG(d, a, b, c, x6, S22, (uint8) 0xc040b340);
+		GG2(c, d, a, b, S23, (uint8) 0x265e5a51);
+		GG(b, c, d, a, x0, S24, (uint8) 0xe9b6c7aa);
+		GG(a, b, c, d, x5, S21, (uint8) 0xd62f105d);
+		GG2(d, a, b, c, S22, (uint8) 0x2441453);
+		GG2(c, d, a, b, S23, (uint8) 0xd8a1e681);
+		GG(b, c, d, a, x4, S24, (uint8) 0xe7d3fbc8);
+		GG2(a, b, c, d, S21, (uint8) 0x21e1cde6);
+		GG(d, a, b, c, x14, S22, (uint8) 0xc33707d6);
+		GG(c, d, a, b, x3, S23, (uint8) 0xf4d50d87);
+		GG2(b, c, d, a, S24, (uint8) 0x455a14ed);
+		GG2(a, b, c, d, S21, (uint8) 0xa9e3e905);
+		GG(d, a, b, c, x2, S22, (uint8) 0xfcefa3f8);
+		GG(c, d, a, b, x7, S23, (uint8) 0x676f02d9);
+		GG2(b, c, d, a, S24, (uint8) 0x8d2a4c8a);
+
+		HH(a, b, c, d, x5, S31, (uint8) 0xfffa3942);
+		HH2(d, a, b, c, S32, (uint8) 0x8771f681);
+		HH2(c, d, a, b, S33, (uint8) 0x6d9d6122);
+		HH(b, c, d, a, x14, S34, (uint8) 0xfde5380c);
+		HH(a, b, c, d, x1, S31, (uint8) 0xa4beea44);
+		HH(d, a, b, c, x4, S32, (uint8) 0x4bdecfa9);
+		HH(c, d, a, b, x7, S33, (uint8) 0xf6bb4b60);
+		HH2(b, c, d, a, S34, (uint8) 0xbebfbc70);
+		HH2(a, b, c, d, S31, (uint8) 0x289b7ec6);
+		HH(d, a, b, c, x0, S32, (uint8) 0xeaa127fa);
+		HH(c, d, a, b, x3, S33, (uint8) 0xd4ef3085);
+		HH(b, c, d, a, x6, S34, (uint8) 0x4881d05);
+		HH2(a, b, c, d, S31, (uint8) 0xd9d4d039);
+		HH2(d, a, b, c, S32, (uint8) 0xe6db99e5);
+		HH2(c, d, a, b, S33, (uint8) 0x1fa27cf8);
+		HH(b, c, d, a, x2, S34, (uint8) 0xc4ac5665);
+
+		II(a, b, c, d, x0, S41, (uint8) 0xf4292244);
+		II(d, a, b, c, x7, S42, (uint8) 0x432aff97);
+		II(c, d, a, b, x14, S43, (uint8) 0xab9423a7);
+		II(b, c, d, a, x5, S44, (uint8) 0xfc93a039);
+		II2(a, b, c, d, S41, (uint8) 0x655b59c3);
+		II(d, a, b, c, x3, S42, (uint8) 0x8f0ccc92);
+		II2(c, d, a, b, S43, (uint8) 0xffeff47d);
+		II(b, c, d, a, x1, S44, (uint8) 0x85845dd1);
+		II2(a, b, c, d, S41, (uint8) 0x6fa87e4f);
+		II2(d, a, b, c, S42, (uint8) 0xfe2ce6e0);
+		II(c, d, a, b, x6, S43, (uint8) 0xa3014314);
+		II2(b, c, d, a, S44, (uint8) 0x4e0811a1);
+		II(a, b, c, d, x4, S41, (uint8) 0xf7537e82);
+		II2(d, a, b, c, S42, (uint8) 0xbd3af235);
+		II(c, d, a, b, x2, S43, (uint8) 0x2ad7d2bb);
+		II2(b, c, d, a, S44, (uint8) 0xeb86d391);
+
+		x0 = a + (uint8) 0x67452301;
+		x1 = b + (uint8) 0xefcdab89;
+		x2 = c + (uint8) 0x98badcfe;
+		x3 = d + (uint8) 0x10325476;
 	} while (--count);
 
-	data_out[idx].v[0] = x0;
-	data_out[idx].v[1] = x1;
-	data_out[idx].v[2] = x2;
-	data_out[idx].v[3] = x3;
-}
\ No newline at end of file
+	char cracked;
+#define K3(q)\
+	cracked = 1;\
+	cracked &= (x0.s##q == salt->hash[0]);\
+	cracked &= (x1.s##q == salt->hash[1]);\
+	cracked &= (x2.s##q == salt->hash[2]);\
+	cracked &= (x3.s##q == salt->hash[3]);\
+	data_out[idx*8+q].cracked = cracked;
+
+
+	K3(0)
+	    K3(1)
+	    K3(2)
+	    K3(3)
+	    K3(4)
+	    K3(5)
+	    K3(6)
+	    K3(7)
+}
diff -urpN magnum-jumbo/src/opencl_phpass_fmt.c magnum-jumbo_phpass-ocl//src/opencl_phpass_fmt.c
--- magnum-jumbo/src/opencl_phpass_fmt.c	2012-08-21 17:35:10.089759332 +0000
+++ magnum-jumbo_phpass-ocl//src/opencl_phpass_fmt.c	2012-08-21 17:39:37.000000000 +0000
@@ -1,5 +1,5 @@
 /*
-* This software is Copyright (c) 2011-2012 Lukas Odzioba <ukasz@openwall.net>
+* This software is Copyright (c) 2011-2012 Lukas Odzioba <ukasz at openwall dot net> 
 * and it is hereby released to the general public under the following terms:
 * Redistribution and use in source and binary forms, with or without modification, are permitted.
 */
@@ -12,26 +12,25 @@
 
 #include "common-opencl.h"
 
-#define uint32_t		unsigned int
-#define uint8_t			unsigned char
+#define uint32_t                unsigned int
+#define uint8_t                 unsigned char
 
-#define FORMAT_LABEL		"phpass-opencl"
-#define FORMAT_NAME		"phpass MD5"
+#define FORMAT_LABEL            "phpass-opencl"
+#define FORMAT_NAME             "phpass MD5"
 
-#define ALGORITHM_NAME		"OpenCL"
+#define ALGORITHM_NAME          "OpenCL"
 
-#define BENCHMARK_COMMENT	" ($P$9 length 8)"
-#define BENCHMARK_LENGTH	-1
+#define BENCHMARK_COMMENT       " ($P$9 length 8)"
+#define BENCHMARK_LENGTH        -1
 
-#define PLAINTEXT_LENGTH	15
-#define CIPHERTEXT_LENGTH	34	/// header = 3 | loopcnt = 1 | salt = 8 | ciphertext = 22
-#define BINARY_SIZE		16
-#define SALT_SIZE		8
-
-#define KEYS_PER_CRYPT		1024*9
-#define MIN_KEYS_PER_CRYPT	KEYS_PER_CRYPT
-#define MAX_KEYS_PER_CRYPT	KEYS_PER_CRYPT
+#define PLAINTEXT_LENGTH        15
+#define CIPHERTEXT_LENGTH       34	/// header = 3 | loopcnt = 1 | salt = 8 | ciphertext = 22
+#define BINARY_SIZE            16
+#define SALT_SIZE               sizeof(phpass_salt)
 
+#define KEYS_PER_CRYPT          64*8*128	//1024*8*40
+#define MIN_KEYS_PER_CRYPT      KEYS_PER_CRYPT
+#define MAX_KEYS_PER_CRYPT      KEYS_PER_CRYPT
 
 //#define _PHPASS_DEBUG
 
@@ -41,54 +40,63 @@ typedef struct {
 } phpass_password;
 
 typedef struct {
-	uint32_t v[4];		///128bits for hash
-} phpass_hash;
+	uint8_t salt[8];
+	uint32_t hash[4];
+	uint32_t rounds;
+} phpass_salt;
+
+typedef struct {
+	uint8_t cracked;
+} phpass_crack;
 
-static phpass_password *inbuffer;//[MAX_KEYS_PER_CRYPT];			/** plaintext ciphertexts **/
-static phpass_hash *outbuffer;//[MAX_KEYS_PER_CRYPT];			/** calculated hashes **/
-static const char phpass_prefix[] = "$P$";
-static char currentsalt[SALT_SIZE + 1];
+
+static phpass_password *inbuffer;			/** plaintext ciphertexts **/
+static phpass_crack *outbuffer;			/** calculated hashes **/
+static phpass_salt *currentsalt;
+static const char phpassP_prefix[] = "$P$";
+static const char phpassH_prefix[] = "$H$";
 
 extern void mem_init(unsigned char *, uint32_t *, char *, char *, int);
 extern void mem_clear(void);
 extern void gpu_phpass(void);
 
 // OpenCL variables:
-static cl_mem mem_in, mem_out, mem_setting;
+static cl_mem mem_in, mem_out, mem_salt;
 static size_t insize = sizeof(phpass_password) * KEYS_PER_CRYPT;
-static size_t outsize = sizeof(phpass_hash) * KEYS_PER_CRYPT;
-static size_t settingsize = sizeof(uint8_t) * SALT_SIZE + 4;
-
+static size_t outsize = sizeof(phpass_crack) * KEYS_PER_CRYPT;
+static size_t saltsize = sizeof(phpass_salt);
 
 static struct fmt_tests tests[] = {
-	/*{"$P$900000000jPBDh/JWJIyrF0.DmP7kT.", "ala"},
-	   {"$P$900000000a94rg7R/nUK0icmALICKj1", "john"},
-	   {"$P$900000001ahWiA6cMRZxkgUxj4x/In0", "john"},
-	   {"$P$900000000m6YEJzWtTmNBBL4jypbHv1", "openwall"},
-	   {"$P$900000000zgzuX4Dc2091D8kak8RdR0", "h3ll00"},
-	   {"$P$900000000qZTL5A0XQUX9hq0t8SoKE0", "1234567890"},
-	   {"$P$900112200B9LMtPy2FSq910c1a6BrH0", "1234567890"},
-	   {"$P$900000000a94rg7R/nUK0icmALICKj1", "john"},
-	   {"$P$9sadli2.wzQIuzsR2nYVhUSlHNKgG/0", "john"},
-	   {"$P$90000000000tbNYOc9TwXvLEI62rPt1", ""}, */
-
-	/*{"$P$9saltstriAcRMGl.91RgbAD6WSq64z.", "a"},
-	   {"$P$9saltstriMljTzvdluiefEfDeGGQEl/", "ab"},
+//      {"$P$900000000jPBDh/JWJIyrF0.DmP7kT.", "ala"},
+//         {"$P$900000000a94rg7R/nUK0icmALICKj1", "john"},
+//         {"$P$900000001ahWiA6cMRZxkgUxj4x/In0", "john"},
+//         {"$P$900000000m6YEJzWtTmNBBL4jypbHv1", "openwall"},
+/*         {"$P$900000000zgzuX4Dc2091D8kak8RdR0", "h3ll00"},
+            {"$P$900000000qZTL5A0XQUX9hq0t8SoKE0", "1234567890"},
+          {"$P$900112200B9LMtPy2FSq910c1a6BrH0", "1234567890"},
+           {"$P$900000000a94rg7R/nUK0icmALICKj1", "john"},
+           {"$P$9sadli2.wzQIuzsR2nYVhUSlHNKgG/0", "john"},
+           {"$P$90000000000tbNYOc9TwXvLEI62rPt1", ""}, 
+*/
+	{"$P$9saltstriAcRMGl.91RgbAD6WSq64z.", "a"},
+	/*   {"$P$9saltstriMljTzvdluiefEfDeGGQEl/", "ab"},
 	   {"$P$9saltstrikCftjZCE7EY2Kg/pjbl8S.", "abc"},
 	   {"$P$9saltstriV/GXRIRi9UVeMLMph9BxF0", "abcd"},
 	   {"$P$9saltstri3JPgLni16rBZtI03oeqT.0", "abcde"},
 	   {"$P$9saltstri0D3A6JyITCuY72ZoXdejV.", "abcdef"},
-	   {"$P$9saltstriXeNc.xV8N.K9cTs/XEn13.", "abcdefg"}, */
-	{"$P$9saltstrinwvfzVRP3u1gxG2gTLWqv.", "abcdefgh"},
-	/*
-	   {"$P$9saltstriSUQTD.yC2WigjF8RU0Q.Z.", "abcdefghi"},
+	   {"$P$9saltstriXeNc.xV8N.K9cTs/XEn13.", "abcdefg"},
+	   {"$P$9saltstrinwvfzVRP3u1gxG2gTLWqv.", "abcdefgh"},
+
+	   {"$H$9saltstriSUQTD.yC2WigjF8RU0Q.Z.", "abcdefghi"},
 	   {"$P$9saltstriWPpGLG.jwJkwGRwdKNEsg.", "abcdefghij"},
 	   {"$P$9saltstrizjDEWUMXTlQHQ3/jhpR4C.", "abcdefghijk"},
 	   {"$P$9saltstriGLUwnE6bl91BPJP6sxyka.", "abcdefghijkl"},
 	   {"$P$9saltstriq7s97e2m7dXnTEx2mtPzx.", "abcdefghijklm"},
 	   {"$P$9saltstriTWMzWKsEeiE7CKOVVU.rS0", "abcdefghijklmn"},
-	   {"$P$9saltstriXt7EDPKtkyRVOqcqEW5UU.", "abcdefghijklmno"},
+	   {"$P$9saltstriXt7EDPKtkyRVOqcqEW5UU.", "abcdefghijklmno"}, 
+	   {"$P$9RjH.g0cuFtd6TnI/A5MRR90TXPc43/", "password__1"},
 	 */
+
 	{NULL}
 };
 
@@ -96,19 +104,20 @@ static void release_all(void)
 {
 	HANDLE_CLERROR(clReleaseKernel(crypt_kernel), "Release Kernel");
 	HANDLE_CLERROR(clReleaseMemObject(mem_in), "Release mem in");
-	HANDLE_CLERROR(clReleaseMemObject(mem_setting), "Release mem setting");
+	HANDLE_CLERROR(clReleaseMemObject(mem_salt), "Release mem setting");
 	HANDLE_CLERROR(clReleaseMemObject(mem_out), "Release mem out");
 	HANDLE_CLERROR(clReleaseCommandQueue(queue[gpu_id]), "Release Queue");
-	MEM_FREE(inbuffer);
-	MEM_FREE(outbuffer);
+	free(inbuffer);
+	free(outbuffer);
 }
 
 static void set_key(char *key, int index)
 {
 #ifdef _PHPASS_DEBUG
-	fprintf(stderr, "set_key(%d) = %s\n", index, key);
+	printf("set_key(%d) = %s\n", index, key);
 #endif
 	int length = strlen(key);
+	memset(inbuffer[index].v, 0, 15);
 	inbuffer[index].length = length;
 	memcpy(inbuffer[index].v, key, length);
 }
@@ -121,14 +130,11 @@ static char *get_key(int index)
 	return ret;
 }
 
-static void init(struct fmt_main *self)
+static void init(struct fmt_main *pFmt)
 {
 	cl_int cl_error;
-
-	global_work_size = MAX_KEYS_PER_CRYPT;
-
-	atexit(release_all);
-	opencl_init("$JOHN/phpass_kernel.cl", gpu_id,platform_id);
+	global_work_size = KEYS_PER_CRYPT / 8;
+	opencl_init("$JOHN/phpass_kernel.cl", gpu_id, platform_id);
 
 	/// Alocate memory
 	inbuffer =
@@ -136,15 +142,17 @@ static void init(struct fmt_main *self)
 	    sizeof(phpass_password));
 	assert(inbuffer != NULL);
 	outbuffer =
-	    (phpass_hash *) calloc(MAX_KEYS_PER_CRYPT,
-	    sizeof(phpass_hash));
+	    (phpass_crack *) calloc(MAX_KEYS_PER_CRYPT, sizeof(phpass_crack));
 	assert(inbuffer != NULL);
+	currentsalt = (phpass_salt *) calloc(1, sizeof(phpass_salt));
+	assert(currentsalt != NULL);
+
 	mem_in =
 	    clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, insize, NULL,
 	    &cl_error);
 	HANDLE_CLERROR(cl_error, "Error alocating mem in");
-	mem_setting =
-	    clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, settingsize,
+	mem_salt =
+	    clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, saltsize,
 	    NULL, &cl_error);
 	HANDLE_CLERROR(cl_error, "Error alocating mem setting");
 	mem_out =
@@ -159,24 +167,30 @@ static void init(struct fmt_main *self)
 		&mem_in), "Error while setting mem_in");
 	HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 1, sizeof(mem_out),
 		&mem_out), "Error while setting mem_out");
-	HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 2, sizeof(mem_setting),
-		&mem_setting), "Error while setting mem_setting");
+	HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 2, sizeof(mem_salt),
+		&mem_salt), "Error while setting mem_salt");
+
+	//opencl_find_best_workgroup(pFmt);
 
-	opencl_find_best_workgroup(self);
+	local_work_size = 64;
+	atexit(release_all);
 }
 
-static int valid(char *ciphertext, struct fmt_main *self)
+
+
+static int valid(char *ciphertext, struct fmt_main *pFmt)
 {
 	uint32_t i, j, count_log2, found;
 
-	int prefix=0;
 	if (strlen(ciphertext) != CIPHERTEXT_LENGTH)
 		return 0;
-	if (strncmp(ciphertext, "$P$", 3) == 0)
-		prefix=1;
-	if (strncmp(ciphertext, "$H$", 3) == 0)
-		prefix=1;
-	if(prefix==0) return 0;
+	found = 0;
+	if (strncmp(ciphertext, phpassP_prefix, 3) == 0)
+		found = 1;
+	if (strncmp(ciphertext, phpassH_prefix, 3) == 0)
+		found = 1;
+	if (!found)
+		return 0;
 
 	for (i = 3; i < CIPHERTEXT_LENGTH; i++) {
 		found = 0;
@@ -195,71 +209,76 @@ static int valid(char *ciphertext, struc
 	return 1;
 };
 
-//code from historical JtR phpass patch
-static void *binary(char *ciphertext)
+
+///code from historical JtR phpass patch
+static void pbinary(char *ciphertext, unsigned char *out)
 {
-	static unsigned char b[BINARY_SIZE];
 	int i, bidx = 0;
 	unsigned sixbits;
 	char *pos = &ciphertext[3 + 1 + 8];
-	memset(b, 0, BINARY_SIZE);
+	memset(out, 0, BINARY_SIZE);
 
 	for (i = 0; i < 5; i++) {
 		sixbits = atoi64[ARCH_INDEX(*pos++)];
-		b[bidx] = sixbits;
+		out[bidx] = sixbits;
 		sixbits = atoi64[ARCH_INDEX(*pos++)];
-		b[bidx++] |= (sixbits << 6);
+		out[bidx++] |= (sixbits << 6);
 		sixbits >>= 2;
-		b[bidx] = sixbits;
+		out[bidx] = sixbits;
 		sixbits = atoi64[ARCH_INDEX(*pos++)];
-		b[bidx++] |= (sixbits << 4);
+		out[bidx++] |= (sixbits << 4);
 		sixbits >>= 4;
-		b[bidx] = sixbits;
+		out[bidx] = sixbits;
 		sixbits = atoi64[ARCH_INDEX(*pos++)];
-		b[bidx++] |= (sixbits << 2);
+		out[bidx++] |= (sixbits << 2);
 	}
 	sixbits = atoi64[ARCH_INDEX(*pos++)];
-	b[bidx] = sixbits;
+	out[bidx] = sixbits;
 	sixbits = atoi64[ARCH_INDEX(*pos++)];
-	b[bidx] |= (sixbits << 6);
+	out[bidx] |= (sixbits << 6);
+}
+
+static void *binary(char *ciphertext)
+{
+	static unsigned char b[BINARY_SIZE];
+	pbinary(ciphertext, b);
 	return (void *) b;
 }
 
 static void *salt(char *ciphertext)
 {
-	static unsigned char salt[SALT_SIZE + 1];
-	memcpy(salt, &ciphertext[4], 8);
-	salt[8] = ciphertext[3];
-	return salt;
+	static phpass_salt salt;
+	salt.rounds = 1 << atoi64[ARCH_INDEX(ciphertext[3])];
+	memcpy(salt.salt, &ciphertext[4], 8);
+	pbinary(ciphertext, (unsigned char *) salt.hash);
+	return &salt;
 }
 
-
 static void set_salt(void *salt)
 {
-	memcpy(currentsalt, salt, SALT_SIZE + 1);
+	memcpy(currentsalt, salt, SALT_SIZE);
 }
 
 static void crypt_all(int count)
 {
 #ifdef _PHPASS_DEBUG
-	fprintf(stderr, "crypt_all(%d)\n", count);
+	printf("crypt_all(%d)\n", count);
 #endif
 	///Prepare setting format: salt+prefix+count_log2
-	char setting[SALT_SIZE + 3 + 1] = { 0 };
-	strcpy(setting, currentsalt);
-	strcpy(setting + SALT_SIZE, phpass_prefix);
-	setting[SALT_SIZE + 3] = atoi64[ARCH_INDEX(currentsalt[8])];
+	int i;
+	memset(outbuffer, 0, outsize);
+
 	/// Copy data to gpu
 	HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_in, CL_FALSE, 0,
 		insize, inbuffer, 0, NULL, NULL), "Copy data to gpu");
-	HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_setting,
-		CL_FALSE, 0, settingsize, setting, 0, NULL, NULL),
+	HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_salt,
+		CL_FALSE, 0, saltsize, currentsalt, 0, NULL, NULL),
 	    "Copy setting to gpu");
 
 	/// Run kernel
 	HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], crypt_kernel, 1,
-		NULL, &global_work_size, &local_work_size, 0, NULL, &profilingEvent),
-	    "Run kernel");
+		NULL, &global_work_size, &local_work_size, 0, NULL,
+		&profilingEvent), "Run kernel");
 	HANDLE_CLERROR(clFinish(queue[gpu_id]), "clFinish");
 
 	/// Read the result back
@@ -268,181 +287,60 @@ static void crypt_all(int count)
 
 	/// Await completion of all the above
 	HANDLE_CLERROR(clFinish(queue[gpu_id]), "clFinish");
-}
-
-static int binary_hash_0(void *binary)
-{
-#ifdef _PHPASS_DEBUG
-	fprintf(stderr, "binary_hash_0 ");
-	int i;
-	uint32_t *b = binary;
-	for (i = 0; i < 4; i++)
-		fprintf(stderr, "%08x ", b[i]);
-	puts("");
-#endif
-	return (((ARCH_WORD_32 *) binary)[0] & 0xf);
-}
-
-static int binary_hash_1(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0xff;
-}
-
-static int binary_hash_2(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0xfff;
-}
-
-static int binary_hash_3(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0xffff;
-}
 
-static int binary_hash_4(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0xfffff;
-}
-
-static int binary_hash_5(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0xffffff;
-}
-
-static int binary_hash_6(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0x7ffffff;
-}
-
-static int get_hash_0(int index)
-{
-#ifdef _PHPASS_DEBUG
-	fprintf(stderr, "get_hash_0:   ");
-	int i;
-	for (i = 0; i < 4; i++)
-		fprintf(stderr, "%08x ", outbuffer[index].v[i]);
-	puts("");
-#endif
-	return outbuffer[index].v[0] & 0xf;
-}
-
-static int get_hash_1(int index)
-{
-	return outbuffer[index].v[0] & 0xff;
-}
-
-static int get_hash_2(int index)
-{
-	return outbuffer[index].v[0] & 0xfff;
-}
-
-static int get_hash_3(int index)
-{
-	return outbuffer[index].v[0] & 0xffff;
-}
-
-static int get_hash_4(int index)
-{
-	return outbuffer[index].v[0] & 0xfffff;
-}
-
-static int get_hash_5(int index)
-{
-	return outbuffer[index].v[0] & 0xffffff;
-}
-
-static int get_hash_6(int index)
-{
-	return outbuffer[index].v[0] & 0x7ffffff;
 }
 
 static int cmp_all(void *binary, int count)
 {
-
-	uint32_t b = ((uint32_t *) binary)[0];
-	uint32_t i;
+	int i, any_cracked = 0;
 	for (i = 0; i < count; i++) {
-		if (b == outbuffer[i].v[0]) {
-#ifdef _PHPASS_DEBUG
-			puts("cmp_all = 1");
-#endif
-			return 1;
-		}
+		any_cracked |= outbuffer[i].cracked;
 	}
-#ifdef _PHPASS_DEBUG
-	puts("cmp_all = 0");
-#endif	/* _PHPASS_DEBUG */
-	return 0;
+	return any_cracked;
 }
 
 static int cmp_one(void *binary, int index)
 {
-	int i;
-	uint32_t *t = (uint32_t *) binary;
-	for (i = 0; i < 4; i++)
-		if (t[i] != outbuffer[index].v[i]) {
-#ifdef _PHPASS_DEBUG
-			puts("cmp_one = 0");
-#endif
-			return 0;
-		}
-#ifdef _PHPASS_DEBUG
-	puts("cmp_one = 1");
-#endif
-	return 1;
+	return outbuffer[index].cracked;
 }
 
-static int cmp_exact(char *source, int count)
+static int cmp_exact(char *source, int index)
 {
-	return 1;
+	return outbuffer[index].cracked;
 }
 
 struct fmt_main fmt_opencl_phpass = {
 	{
-		FORMAT_LABEL,
-		FORMAT_NAME,
-		ALGORITHM_NAME,
-		BENCHMARK_COMMENT,
-		BENCHMARK_LENGTH,
-		PLAINTEXT_LENGTH,
-		BINARY_SIZE,
-		SALT_SIZE + 1,
-		MIN_KEYS_PER_CRYPT,
-		MAX_KEYS_PER_CRYPT,
-		FMT_CASE | FMT_8_BIT,
-		tests
-	}, {
-		init,
-		fmt_default_prepare,
-		valid,
-		fmt_default_split,
-		binary,
-		salt,
-		{
-			binary_hash_0,
-			binary_hash_1,
-			binary_hash_2,
-			binary_hash_3,
-			binary_hash_4,
-			binary_hash_5,
-			binary_hash_6
-		},
-		fmt_default_salt_hash,
-		set_salt,
-		set_key,
-		get_key,
-		fmt_default_clear_keys,
-		crypt_all,
-		{
-			get_hash_0,
-			get_hash_1,
-			get_hash_2,
-			get_hash_3,
-			get_hash_4,
-			get_hash_5,
-			get_hash_6
-		},
-		cmp_all,
-		cmp_one,
-		cmp_exact
-	}
+		    FORMAT_LABEL,
+		    FORMAT_NAME,
+		    ALGORITHM_NAME,
+		    BENCHMARK_COMMENT,
+		    BENCHMARK_LENGTH,
+		    PLAINTEXT_LENGTH,
+		    BINARY_SIZE,
+		    SALT_SIZE,
+		    MIN_KEYS_PER_CRYPT,
+		    MAX_KEYS_PER_CRYPT,
+		    FMT_CASE | FMT_8_BIT,
+	    tests},
+	{
+		    init,
+		    fmt_default_prepare,
+		    valid,
+		    fmt_default_split,
+		    binary,
+		    salt,
+		    {
+			fmt_default_binary_hash},
+		    fmt_default_salt_hash,
+		    set_salt,
+		    set_key,
+		    get_key,
+		    fmt_default_clear_keys,
+		    crypt_all,
+		    {
+			fmt_default_get_hash},
+		    cmp_all,
+		    cmp_one,
+	    cmp_exact}
 };