diff -urpN magnum-jumbo/src/opencl/phpass_kernel.cl magnum-jumbo_phpass_oldtest//src/opencl/phpass_kernel.cl
--- magnum-jumbo/src/opencl/phpass_kernel.cl	2012-06-25 00:40:42.738451690 +0000
+++ magnum-jumbo_phpass_oldtest//src/opencl/phpass_kernel.cl	2012-06-25 00:50:52.014452056 +0000
@@ -1,65 +1,65 @@
 /*
-* This software is Copyright (c) 2011,2012 Lukas Odzioba <ukasz@openwall.com> 
+* 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 SALT_SIZE		8
-#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : disable
-
 typedef struct {
-	//uint v[4]; //15 bytes password + last byte for length
-	uchar v[PLAINTEXT_LENGTH];
-	uchar length;
+	unsigned char v[PLAINTEXT_LENGTH];
+	unsigned char length;
 } phpass_password;
 
 typedef struct {
-	uint v[4];
+	unsigned int v[4];
 } phpass_hash;
 
-#define ROTATE_LEFT(x, s)	rotate(x,s)
-//#define F(x, y, z)		((z) ^ ((x) & ((y) ^ (z))))
-//#define G(x, y, z)		((y) ^ ((z) & ((x) ^ (y))))
-
-#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 ROTATE_LEFT(x, s) ((x << s) | (x >> (32 - s)))
+#define F(x, y, z) (((x) & (y)) | ((~x) & (z)))
+#define G(x, y, z) (((x) & (z)) | ((y) & (~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) + (ac); \
-   (a) = ROTATE_LEFT ((a), (s)) + (b); \
+  {(a) += F ((b), (c), (d)) + (x) + (uint32_t)(ac); \
+   (a) = ROTATE_LEFT ((a), (s)); \
+   (a) += (b); \
   }
 #define GG(a, b, c, d, x, s, ac) \
-  {(a) += G ((b), (c), (d)) + (x) + (ac); \
-   (a) = ROTATE_LEFT ((a), (s)) + (b); \
+  {(a) += G ((b), (c), (d)) + (x) + (uint32_t)(ac); \
+   (a) = ROTATE_LEFT ((a), (s)); \
+   (a) += (b); \
   }
 #define HH(a, b, c, d, x, s, ac) \
-  {(a) += H ((b), (c), (d)) + (x) + (ac); \
-   (a) = ROTATE_LEFT ((a), (s)) + (b); \
+  {(a) += H ((b), (c), (d)) + (x) + (uint32_t)(ac); \
+   (a) = ROTATE_LEFT ((a), (s)); \
+   (a) += (b); \
   }
 #define II(a, b, c, d, x, s, ac) \
-  {(a) += I ((b), (c), (d)) + (x) + (ac); \
-   (a) = ROTATE_LEFT ((a), (s)) + (b); \
+  {(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 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
@@ -69,27 +69,20 @@ typedef struct {
 
 
 
-inline void md5(uint4 len,__private uint4 * internal_ret,__private uint4 * x)
+inline void cuda_md5(char len,__private uint32_t * internal_ret,__private uint32_t * x)
 {
-	uint4 x14 = len << 3;
+	x[len / 4] |= (((uint32_t) 0x80) << ((len & 0x3) << 3));
+	uint32_t x14 = len << 3;
 
-	uint4 a;
-	uint4 b = 0xefcdab89;
-	uint4 c = 0x98badcfe;
-	uint4 d = 0x10325476;
-
-	a = AC1 + x[0];
-	a = ROTATE_LEFT(a, S11);
-	a += b;			/* 1 */
-	d = (c ^ (a & MASK1)) + x[1] + AC2pCd;
-	d = ROTATE_LEFT(d, S12);
-	d += a;			/* 2 */
-	c = F(d, a, b) + x[2] + AC3pCc;
-	c = ROTATE_LEFT(c, S13);
-	c += d;			/* 3 */
-	b = F(c, d, a) + x[3] + AC4pCb;
-	b = ROTATE_LEFT(b, S14);
-	b += c;
+	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);
@@ -160,132 +153,150 @@ inline void md5(uint4 len,__private uint
 	internal_ret[3] = d + 0x10325476;
 }
 
-inline void clear_ctx(__private uint4 * x)
+inline void clear_ctx(__private uint32_t * x)
 {
-	uint4 zero = (uint4) (0, 0,0,0);
-	for (int i = 0; i < 8; i++)
-		x[i] = zero;
-}
-inline void clean_ctx(__private uint *x){
-	for(int i=0;i<8;i++)
-		x[i]=0;
+	int i;
+	for (i = 0; i < 8; i++)
+		*x++ = 0;
 }
 
 
+
 __kernel void phpass
     (   __global    const   phpass_password*    data
     ,   __global            phpass_hash*    	data_out
     ,   __global    const   char* 		setting
     )
 {
-	uint4 x[8],length;
-	uint sx[8],i,idx = get_global_id(0);
-	uint count = 1 << setting[SALT_SIZE+3];
-
+	uint32_t x[8];
 	clear_ctx(x);
+	uint32_t a, b, c, d, x0, x1, x2, x3, x4, x5, x6, x7;
 
-	__global const uchar *password0=data[idx*4+0].v;
-	__global const uchar *password1=data[idx*4+1].v;
-	__global const uchar *password2=data[idx*4+2].v;
-	__global const uchar *password3=data[idx*4+3].v;
-
-
-
-	length.s0=data[idx*4+0].length;
-	length.s1=data[idx*4+1].length;
-	length.s2=data[idx*4+2].length;
-	length.s3=data[idx*4+3].length;
-
-	__private uint *buff2=(uint*)sx;
-	#define K1(q)\
-		clean_ctx(sx);\
-		buff2[0]=(setting[3]<<24)|setting[2]<<16|setting[1]<<8|setting[0];\
-		buff2[1]=(setting[7]<<24)|setting[6]<<16|setting[5]<<8|setting[4];\
-		buff2[2]=(password##q[3]<<24)|password##q[2]<<16|password##q[1]<<8|password##q[0];\
-		buff2[3]=(password##q[7]<<24)|password##q[6]<<16|password##q[5]<<8|password##q[4];\
-		buff2[4]=(password##q[11]<<24)|password##q[10]<<16|password##q[9]<<8|password##q[8];\
-		buff2[5]=password##q[14]<<16|password##q[13]<<8|password##q[12];\
-		for ( i = 0; i < 8; i++)\
-			x[i].s##q=sx[i];
-		K1(0);
-		K1(1);
-		K1(2);
-		K1(3);
-	#undef K1 
-
-
-/*
- 	#define K1(q)\
-		clean_ctx(sx);\
-		for (i = 0; i < 8; i++)\
-			buff[i] = setting[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);
-	#undef K1 
-*/
-
-	uint4 len=length+(uint4)(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));
-
-	md5(len, x, x);
-
-
-
-/*#define K2(q)\
-		clean_ctx(sx);\
-		for(i=0;i<length.s##q;i++)\
-			buff[i]=password##q[i];\
-		for(i=0;i<2;i++)\
-			x[i+4].s##q=sx[i];
-	K2(0);
-	K2(1);
-	K2(2);
-	K2(3);
-#undef K2
-*/
+	uint32_t idx = get_global_id(0);
 
-#define K2(q)\
-		clean_ctx(sx);\
-		buff2[0]=(password##q[3]<<24)|password##q[2]<<16|password##q[1]<<8|password##q[0];\
-		buff2[1]=(password##q[7]<<24)|password##q[6]<<16|password##q[5]<<8|password##q[4];\
-		buff2[2]=(password##q[11]<<24)|password##q[10]<<16|password##q[9]<<8|password##q[8];\
-		buff2[3]=password##q[14]<<16|password##q[13]<<8|password##q[12];\
-		for(i=0;i<2;i++)\
-			x[i+4].s##q=sx[i];
-	K2(0);
-	K2(1);
-	K2(2);
-	K2(3);
-#undef K2
-
-	len = 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));
+	__global const char *password = (__global const char*) data[idx].v;
+	int length, count, i;
+	__private unsigned char *buff = (unsigned char *) x;
+
+	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];
+	}
+
+	cuda_md5(8 + length, x, x);
+	count = 1 << setting[SALT_SIZE+3];
+	for (i = 16; i < 16 + length; i++)
+		buff[i] = password[i - 16];
+
+
+	uint32_t len = 16 + length;
+	uint32_t x14 = len << 3;
+
+	x[len / 4] |= ((0x80) << ((len & 0x3) << 3));
+	x0 = x[0];
+	x1 = x[1];
+	x2 = x[2];
+	x3 = x[3];
+	x4 = x[4];
+	x5 = x[5];
+	x6 = x[6];
+	x7 = x[7];
+do {
+
+		b = 0xefcdab89;
+		c = 0x98badcfe;
+		d = 0x10325476;
+
+// FF(a, b, c, d, x0, S11, 0xd76aa478);
+		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, 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;
 
-	do {
-	  md5(len, x, x);
 	} while (--count);
 
-#define K3(q)\
-	data_out[idx*4+q].v[0]=x[0].s##q;\
-	data_out[idx*4+q].v[1]=x[1].s##q;\
-	data_out[idx*4+q].v[2]=x[2].s##q;\
-	data_out[idx*4+q].v[3]=x[3].s##q;
-	K3(0)
-	K3(1)
-	K3(2)
-	K3(3)
-#undef K3
+	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
diff -urpN magnum-jumbo/src/opencl_phpass_fmt.c magnum-jumbo_phpass_oldtest//src/opencl_phpass_fmt.c
--- magnum-jumbo/src/opencl_phpass_fmt.c	2012-06-25 00:40:42.739451675 +0000
+++ magnum-jumbo_phpass_oldtest//src/opencl_phpass_fmt.c	2012-06-25 01:02:22.369576850 +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.
 */
@@ -15,22 +15,24 @@
 #define uint32_t		unsigned int
 #define uint8_t			unsigned char
 
-#define FORMAT_LABEL		"phpass-opencl"
-#define FORMAT_NAME		"phpass MD5"
-
-#define ALGORITHM_NAME		"OpenCL"
-
-#define BENCHMARK_COMMENT	" ($P$9 length 8)"
-#define BENCHMARK_LENGTH	-1
+#define PHPASS_TYPE		"PORTABLE-MD5"
 
 #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*4
+#define KEYS_PER_CRYPT		1024*9
 #define MIN_KEYS_PER_CRYPT	KEYS_PER_CRYPT
 #define MAX_KEYS_PER_CRYPT	KEYS_PER_CRYPT
+#define FORMAT_LABEL		"phpass-opencl"
+#define FORMAT_NAME		"phpass MD5"
+
+#define ALGORITHM_NAME		"OpenCL"
+
+#define BENCHMARK_COMMENT	" ($P$9 length 8)"
+#define BENCHMARK_LENGTH	-1
+
 
 //#define _PHPASS_DEBUG
 
@@ -43,8 +45,8 @@ typedef struct {
 	uint32_t v[4];		///128bits for hash
 } phpass_hash;
 
-static phpass_password *inbuffer;//[MAX_KEYS_PER_CRYPT];			/** plaintext ciphertexts **/
-static phpass_hash *outbuffer;//[MAX_KEYS_PER_CRYPT];			/** calculated hashes **/
+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];
 
@@ -57,7 +59,7 @@ static cl_mem mem_in, mem_out, mem_setti
 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 global_work_size = KEYS_PER_CRYPT/4;
+static size_t global_work_size = KEYS_PER_CRYPT;
 
 
 static struct fmt_tests tests[] = {
@@ -72,15 +74,15 @@ static struct fmt_tests tests[] = {
 	   {"$P$9sadli2.wzQIuzsR2nYVhUSlHNKgG/0", "john"},
 	   {"$P$90000000000tbNYOc9TwXvLEI62rPt1", ""}, */
 
-	/*{"$P$9saltstriAcRMGl.91RgbAD6WSq64z.", "a"},
+	{"$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$9saltstriXeNc.xV8N.K9cTs/XEn13.", "abcdefg"}, 
 	{"$P$9saltstrinwvfzVRP3u1gxG2gTLWqv.", "abcdefgh"},
-	/*
+	
 	   {"$P$9saltstriSUQTD.yC2WigjF8RU0Q.Z.", "abcdefghi"},
 	   {"$P$9saltstriWPpGLG.jwJkwGRwdKNEsg.", "abcdefghij"},
 	   {"$P$9saltstrizjDEWUMXTlQHQ3/jhpR4C.", "abcdefghijk"},
@@ -88,7 +90,7 @@ static struct fmt_tests tests[] = {
 	   {"$P$9saltstriq7s97e2m7dXnTEx2mtPzx.", "abcdefghijklm"},
 	   {"$P$9saltstriTWMzWKsEeiE7CKOVVU.rS0", "abcdefghijklmn"},
 	   {"$P$9saltstriXt7EDPKtkyRVOqcqEW5UU.", "abcdefghijklmno"}, 
-	 */
+	 
 	{NULL}
 };
 
@@ -99,8 +101,6 @@ static void release_all(void)
 	HANDLE_CLERROR(clReleaseMemObject(mem_setting), "Release mem setting");
 	HANDLE_CLERROR(clReleaseMemObject(mem_out), "Release mem out");
 	HANDLE_CLERROR(clReleaseCommandQueue(queue[gpu_id]), "Release Queue");
-	free(inbuffer);
-	free(outbuffer);
 }
 
 static void set_key(char *key, int index)
@@ -121,21 +121,71 @@ static char *get_key(int index)
 	return ret;
 }
 
+static void find_best_workgroup()
+{
+	cl_event myEvent;
+	cl_ulong startTime, endTime, kernelExecTimeNs = CL_ULONG_MAX;
+	size_t my_work_group = 1;
+	cl_int ret_code;
+	int i;
+	size_t max_group_size;
+	clGetDeviceInfo(devices[gpu_id], CL_DEVICE_MAX_WORK_GROUP_SIZE,
+	    sizeof(max_group_size), &max_group_size, NULL);
+	cl_command_queue queue_prof =
+	    clCreateCommandQueue(context[gpu_id], devices[gpu_id],
+	    CL_QUEUE_PROFILING_ENABLE,
+	    &ret_code);
+	HANDLE_CLERROR(ret_code, "Error while creating command queue");
+	local_work_size = 1;
+	/// Set keys
+	char *pass = "aaaaaaaa";
+	for (i = 0; i < KEYS_PER_CRYPT; i++) {
+		set_key(pass, i);
+	}
+	///Set salt
+	memcpy(currentsalt, "saltstri9", 9);
+	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])];
+
+	///Copy data to GPU
+	HANDLE_CLERROR(clEnqueueWriteBuffer(queue_prof, mem_in, CL_FALSE, 0,
+		insize, inbuffer, 0, NULL, NULL), "Copy data to gpu");
+	HANDLE_CLERROR(clEnqueueWriteBuffer(queue_prof, mem_setting, CL_FALSE,
+		0, settingsize, setting, 0, NULL, NULL),
+	    "Copy setting to gpu");
+
+	///Find best local work size
+	for (my_work_group = 1; (int) my_work_group <= (int) max_group_size;
+	    my_work_group *= 2) {
+
+		HANDLE_CLERROR(clEnqueueNDRangeKernel(queue_prof, crypt_kernel,
+			1, NULL, &global_work_size, &my_work_group, 0, NULL,
+			&myEvent), "Run kernel");
+
+		HANDLE_CLERROR(clFinish(queue_prof), "clFinish error");
+		clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_SUBMIT,
+		    sizeof(cl_ulong), &startTime, NULL);
+		clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_END,
+		    sizeof(cl_ulong), &endTime, NULL);
+
+		if ((endTime - startTime) < kernelExecTimeNs) {
+			kernelExecTimeNs = endTime - startTime;
+			local_work_size = my_work_group;
+		}
+		//printf("%d time=%lld\n",(int) my_work_group, endTime-startTime);
+	}
+	printf("Optimal local work size = %d\n", (int) local_work_size);
+	clReleaseCommandQueue(queue_prof);
+}
+
 static void init(struct fmt_main *pFmt)
 {
-	atexit(release_all);
+	//atexit(release_all);
 	opencl_init("$JOHN/phpass_kernel.cl", gpu_id,platform_id);
 
 	/// Alocate memory
-	inbuffer =
-	    (phpass_password *) calloc(MAX_KEYS_PER_CRYPT,
-	    sizeof(phpass_password));
-	assert(inbuffer != NULL);
-	outbuffer =
-	    (phpass_hash *) calloc(MAX_KEYS_PER_CRYPT,
-	    sizeof(phpass_hash));
-	assert(inbuffer != NULL);
-
 	cl_int cl_error;
 	mem_in =
 	    clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, insize, NULL,
@@ -153,14 +203,11 @@ static void init(struct fmt_main *pFmt)
 	/// Setup kernel parameters
 	crypt_kernel = clCreateKernel(program[gpu_id], "phpass", &cl_error);
 	HANDLE_CLERROR(cl_error, "Error creating kernel");
-	HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 0, sizeof(mem_in),
-		&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");
+	clSetKernelArg(crypt_kernel, 0, sizeof(mem_in), &mem_in);
+	clSetKernelArg(crypt_kernel, 1, sizeof(mem_out), &mem_out);
+	clSetKernelArg(crypt_kernel, 2, sizeof(mem_setting), &mem_setting);
 
-	opencl_find_best_workgroup(pFmt);
+	find_best_workgroup();
 }
 
 static int valid(char *ciphertext, struct fmt_main *pFmt)
@@ -252,7 +299,7 @@ static void crypt_all(int count)
 
 	/// Run kernel
 	HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], crypt_kernel, 1,
-		NULL, &global_work_size, &local_work_size, 0, NULL, &profilingEvent),
+		NULL, &global_work_size, &local_work_size, 0, NULL, NULL),
 	    "Run kernel");
 	HANDLE_CLERROR(clFinish(queue[gpu_id]), "clFinish");