diff --git a/run/dumb16.conf b/run/dumb16.conf
index bdb3d73..791fc41 100644
--- a/run/dumb16.conf
+++ b/run/dumb16.conf
@@ -25,9 +25,6 @@ void init()
 	minlength = 1;	// Initial password length to try, must be at least 1
 	maxlength = 2;	// Must be at least same as minlength
 
-	# UTF-8 representation is up to 3x the character length
-	if (maxlength * 3 > cipher_limit)
-		maxlength = (cipher_limit + 2) / 3;
 /*
  * This defines the character set. This is auto-generated from UnicodeData.txt
  * and we skip control characters.
diff --git a/run/dumb32.conf b/run/dumb32.conf
index 16071ab..725e716 100644
--- a/run/dumb32.conf
+++ b/run/dumb32.conf
@@ -25,9 +25,6 @@ void init()
 	minlength = 1;	// Initial password length to try, must be at least 1
 	maxlength = 2;	// Must be at least same as minlength
 
-	# UTF-8 representation is up to 3x the character length
-	if (maxlength * 3 > cipher_limit)
-		maxlength = (cipher_limit + 2) / 3;
 /*
  * This defines the character set. This is auto-generated from UnicodeData.txt
  * and we skip control characters.
diff --git a/run/pass_gen.pl b/run/pass_gen.pl
index 2c84693..b5b64cc 100755
--- a/run/pass_gen.pl
+++ b/run/pass_gen.pl
@@ -89,10 +89,10 @@ my @i64 = ('.','/','0'..'9','A'..'Z','a'..'z');
 my @ns_i64 = ('A'..'Z', 'a'..'z','0'..'9','+','/',);
 my @userNames = (
 	"admin", "root", "bin", "Joe", "fi15_characters", "Babeface", "Herman", "lexi Conrad", "jack", "John", "sz110",
-	"fR14characters", "Thirteenchars", "Twelve_chars", "elev__chars", "teN__chars", "six16_characters", 
+	"fR14characters", "Thirteenchars", "Twelve_chars", "elev__chars", "teN__chars", "six16_characters",
 #	"Bărtin",
-	"ninechars", "eightchr", "sevench", "barney", "C0ffee", "deadcafe", "user", "01234", "nineteen_characters", 
-	"eight18_characters", "seven17characters", "u1", "harvey", "john", "ripper", "a", "Hank", "1", "u2", "u3", 
+	"ninechars", "eightchr", "sevench", "barney", "C0ffee", "deadcafe", "user", "01234", "nineteen_characters",
+	"eight18_characters", "seven17characters", "u1", "harvey", "john", "ripper", "a", "Hank", "1", "u2", "u3",
 	"2", "3", "usr", "usrx", "usry", "skippy", "Bing", "Johnson", "addams", "anicocls", "twentyXXX_characters",
 	"twentyoneX_characters", "twentytwoXX_characters");
 
@@ -695,7 +695,7 @@ sub moffet_coinflip {
 	my $shift_a; my $shift_b;
 	my $indirect_a; my $indirect_b;
 	my $bit_a; my $bit_b;
-	
+
 	for ($i = 0; $i < 16; $i++) {
 		my $j;
 		$j = ($i + 3) & 0xF;
@@ -724,7 +724,7 @@ sub moffet_coinflip {
 
 	$bit_a = md5bit($c, $indirect_a);
 	$bit_b = md5bit($c, $indirect_b);
-	
+
 	return $bit_a ^ $bit_b;
 }
 sub _sunmd5_hash {
@@ -739,7 +739,7 @@ sub _sunmd5_hash {
 		if (moffet_coinflip($c, $i)) {
 			$c = md5(
 					$c,
-					# this long constant string (AND the null trailing), 
+					# this long constant string (AND the null trailing),
 					# need to be added, then the round's text number
 					"To be, or not to be,--that is the question:--\n",
 					"Whether 'tis nobler in the mind to suffer\n",
@@ -782,7 +782,7 @@ sub _sunmd5_hash {
 		}
 		$i++;
 	}
-	return $c;	
+	return $c;
 }
 sub sunmd5 {
 	if (defined $argsalt) { $salt = $argsalt; } else { $salt=randstr(16); }
@@ -882,16 +882,16 @@ sub _sha_crypts {
 	my $a; my $b, my $c, my $tmp; my $i; my $ds; my $dp; my $p; my $s;
 	my ($func, $bits, $key, $salt) = @_;
 	my $bytes = $bits/8;
-	
+
 	$b = $func->($key.$salt.$key);
-	
+
 	# Add for any character in the key one byte of the alternate sum.
 	$tmp = $key . $salt;
 	for ($i = length($key); $i > 0; $i -= $bytes) {
 		if ($i > $bytes) { $tmp .= $b; }
 		else { $tmp .= substr($b,0,$i); }
 	}
-	
+
 	# Take the binary representation of the length of the key and for every 1 add the alternate sum, for every 0 the key.
 	for ($i = length($key); $i > 0; $i >>= 1) {
 		if (($i & 1) != 0) { $tmp .= $b; }
@@ -899,7 +899,7 @@ sub _sha_crypts {
 	}
 	$a = $func->($tmp);
 	# NOTE, this will be the 'initial' $c value in the inner loop.
-	
+
 	# For every character in the password add the entire password.  produces DP
 	$tmp = "";
 	for ($i = 0; $i < length($key); ++$i) {
@@ -919,13 +919,13 @@ sub _sha_crypts {
 		$tmp .= $salt;
 	}
 	$ds = $func->($tmp);
-	
+
 	# Create byte sequence S.
 	for ($i = length($salt); $i > 0; $i -= $bytes) {
 		if ($i > $bytes) { $s .= $ds; }
 		else { $s .= substr($ds,0,$i); }
 	}
-	
+
 	$c = $a; # Ok, we saved this, which will 'seed' our crypt value here in the loop.
 	# now we do 5000 iterations of md5.
 	for ($i = 0; $i < 5000; ++$i) {
@@ -939,13 +939,13 @@ sub _sha_crypts {
 		$c = $func->($tmp);
 	}
 #	printf ("F =" . unpack("H*", $c) . "\n");  # final value.
-	
-	# $c now contains the 'proper' sha_X_crypt hash.  However, a strange transposition and 
+
+	# $c now contains the 'proper' sha_X_crypt hash.  However, a strange transposition and
 	# base-64 conversion. We do the same here, to get the same hash.  sha256 and sha512 use
 	# a different key schedule.  I have come up with a way to do this, that is not using a
 	# table, but using modular walking of the data, 3 values at a time.
 	# seel http://www.akkadia.org/drepper/SHA-crypt.txt for information
-	
+
 	my $inc1; my $inc2; my $mod; my $end;
 	if ($bits==256) { $inc1=10;$inc2=21;$mod=30;$end=0;  }
 	else            { $inc1=21;$inc2=22;$mod=63;$end=21; }
@@ -971,7 +971,7 @@ sub sha512crypt {
 	print "u$u-sha512crypt:\$6\$$salt\$$bin:$u:0:$_[0]::\n";
 }
 sub xsha512 {
-# simple 4 byte salted crypt.  No seperator char, just raw hash. Also 'may' have $LION$.  We altenate, and every other 
+# simple 4 byte salted crypt.  No seperator char, just raw hash. Also 'may' have $LION$.  We altenate, and every other
 # hash get $LION$ (all even ones)
 	if (defined $argsalt) { $salt = $argsalt; } else { $salt=randstr(4); }
 	print "u$u-XSHA512:";
@@ -1152,7 +1152,7 @@ sub oracle_no_upcase_change {
 	if (!defined $arg_hidden_cp) { print "ERROR, for this format, you MUST use -hiddencp=CP to set the proper code page conversion\n"; exit(1); }
 
 	my $pass = $username . Encode::decode(":".$arg_hidden_cp, $_[0]);
-	
+
 	my $userpass = encode("UTF-16BE", $pass);
 	$userpass .= pack('C', 0) while (length($userpass) % 8);
 	my $key = pack('H*', "0123456789ABCDEF");
@@ -1279,7 +1279,7 @@ sub mschapv2 {
 }
 sub crc_32 {
 	my $pwd = shift;
-	if (rand(256) > 245) { 
+	if (rand(256) > 245) {
 		my $init = rand(2000000000);
 		printf("$u-crc32:\$crc32\$%08x.%08x:0:0:100:%s:\n", $init, crc32($pwd,$init), $pwd);
 	} else {
@@ -1321,7 +1321,7 @@ sub drupal7 {
 	my $h = sha512($salt.$_[0]);
 	my $i = 16384;
 	do { $h = sha512($h.$_[0]); } while (--$i > 0);
-	
+
 	print "u$u-drupal:\$S\$C",$salt,substr(base64i($h),0,43),":$u:0:$_[0]::\n";
 }
 sub epi {
@@ -1346,7 +1346,7 @@ sub nukedclan {
 	my $pass_hash = sha1_hex($_[0]);
 	my $i = 0; my $k;
 	$k = hex($decal);
-	
+
 	my $out = "";
 	for (; $i < 40; $i += 1, $k += 1) {
 		$out .= substr($pass_hash, $i, 1);
@@ -1599,9 +1599,9 @@ sub dynamic_compile {
 			$dynamic_args==100 && do {$fmt='whirlpool($p)';				last SWITCH; };
 			$dynamic_args==101 && do {$fmt='whirlpool($s.$p),saltlen=2';	last SWITCH; };
 			$dynamic_args==102 && do {$fmt='whirlpool($p.$s)';			last SWITCH; };
-			
+
 			# 7, 17, 19, 20, 21, 27, 28 are still handled by 'special' functions.
-			
+
 			# since these are in dynamic.conf, and treatly 'like' builtins, we might as well put them here.
 			$dynamic_args==1001 && do {$fmt='md5(md5(md5(md5($p))))';	last SWITCH; };
 			$dynamic_args==1002 && do {$fmt='md5(md5(md5(md5(md5($p)))))';	last SWITCH; };
@@ -2089,7 +2089,7 @@ sub dynamic_load_salt2() {
 ##########################################################################
 #  Here are the ACTUAL pCode primative functions.  These handle pretty
 # much everything dealing with hashing expressions for md5/md4/sha1/sha224
-# /sha256/sha384/sha512/gost/whirlpool. 
+# /sha256/sha384/sha512/gost/whirlpool.
 # There are some variables which will be properly prepared prior to any of these
 # pCode functions.  These are $gen_pw (the password, possibly in unicode
 # format).  $gen_s (the salt), $gen_s2 (the 2nd salt), $gen_u the username
diff --git a/src/Makefile b/src/Makefile
index e9dcef3..e529d84 100644
--- a/src/Makefile
+++ b/src/Makefile
@@ -115,6 +115,7 @@ JOHN_OBJS = \
 	racf2john.o \
 	pwsafe2john.o \
 	keepass2john.o \
+	keychain2john.o \
 	wpapsk_fmt.o hccap2john.o \
 	mozilla_fmt.o KeyDBCracker.o mozilla_des.o lowpbe.o mozilla2john.o \
 	$(PLUGFORMATS_OBJS) \
@@ -200,13 +201,17 @@ PROJ = ../run/john ../run/unshadow ../run/unafs ../run/unique ../run/undrop \
 	../run/ssh2john ../run/pdf2john ../run/rar2john ../run/zip2john \
 	../run/genmkvpwd ../run/mkvcalcproba ../run/calc_stat \
 	../run/tgtsnarf ../run/racf2john ../run/mozilla2john ../run/hccap2john \
-	../run/pwsafe2john ../run/raw2dyna ../run/keepass2john john.local.conf
+	../run/pwsafe2john ../run/raw2dyna ../run/keepass2john \
+	../run/keychain2john \
+	john.local.conf
 PROJ_DOS = ../run/john.bin ../run/john.com \
 	../run/unshadow.com ../run/unafs.com ../run/unique.com \
 	../run/undrop.com \
 	../run/ssh2john.com ../run/pdf2john.com ../run/rar2john.com ../run/zip2john \
 	../run/racf2john.com ../run/mozilla2john.com ../run/hccap2john.com \
-	../run/pwsafe2john.com ../run/keepass2john.com john.local.conf
+	../run/pwsafe2john.com ../run/keepass2john.com \
+	../run/keychain2john.com \
+	john.local.conf
 PROJ_WIN32 = ../run/john.exe \
 	../run/unshadow.exe ../run/unafs.exe ../run/unique.exe \
 	../run/undrop.exe \
@@ -214,6 +219,7 @@ PROJ_WIN32 = ../run/john.exe \
 	../run/genmkvpwd.exe ../run/mkvcalcproba.exe ../run/calc_stat.exe \
 	../run/racf2john.exe ../run/mozilla2john.exe ../run/hccap2john.exe \
 	../run/pwsafe2john.exe ../run/raw2dyna.exe ../run/keepass2john.exe \
+	../run/keychain2john.exe \
 	john.local.conf
 PROJ_WIN32_MINGW = ../run/john-mingw.exe \
 	../run/unshadow.exe ../run/unafs.exe ../run/unique.exe \
@@ -222,6 +228,7 @@ PROJ_WIN32_MINGW = ../run/john-mingw.exe \
 	../run/genmkvpwd.exe ../run/mkvcalcproba.exe ../run/calc_stat.exe \
 	../run/racf2john.exe ../run/mozilla2john.exe ../run/hccap2john.exe \
 	../run/pwsafe2john.exe ../run/raw2dyna.exe ../run/keepass2john.exe \
+	../run/keychain2john.exe \
 	john.local.conf
 
 default:
@@ -485,7 +492,7 @@ linux-x86-avx:
 	$(LN) x86-sse.h arch.h
 	@echo "#define JOHN_BLD" '"'$@'"' > john_build_rule.h
 	$(MAKE) $(PROJ) \
-		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86.o" \
+		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86.o x86-sse.o sha1-mmx.o md4-mmx.o md5-mmx.o sse-intrinsics.o" \
 		CFLAGS_MAIN="$(CFLAGS) -m32 -DJOHN_AVX -DHAVE_CRYPT" \
 		CFLAGS="$(CFLAGS) -m32 -mavx -DHAVE_CRYPT" \
 		ASFLAGS="$(ASFLAGS) -m32 -mavx" \
@@ -495,7 +502,7 @@ linux-x86-xop:
 	$(LN) x86-sse.h arch.h
 	@echo "#define JOHN_BLD" '"'$@'"' > john_build_rule.h
 	$(MAKE) $(PROJ) \
-		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86.o" \
+		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86.o x86-sse.o sha1-mmx.o md4-mmx.o md5-mmx.o sse-intrinsics.o" \
 		CFLAGS_MAIN="$(CFLAGS) -m32 -DJOHN_XOP -DHAVE_CRYPT" \
 		CFLAGS="$(CFLAGS) -m32 -mxop -DHAVE_CRYPT" \
 		ASFLAGS="$(ASFLAGS) -m32 -mxop" \
@@ -1529,6 +1536,10 @@ cuda_pwsafe_fmt.o: cuda_pwsafe.o cuda_pwsafe_fmt.c
 	$(RM) ../run/keepass2john
 	ln -s john ../run/keepass2john
 
+../run/keychain2john: ../run/john
+	$(RM) ../run/keychain2john
+	ln -s john ../run/keychain2john
+
 ../run/zip2john: ../run/john
 	$(RM) ../run/zip2john
 	ln -s john ../run/zip2john
@@ -1586,6 +1597,9 @@ endif
 ../run/keepass2john.com: john.com
 	copy john.com ..\run\keepass2john.com
 
+../run/keychain2john.com: john.com
+	copy john.com ..\run\keychain2john.com
+
 ../run/pdf2john.com: john.com
 	copy john.com ..\run\pdf2john.com
 
@@ -1658,6 +1672,10 @@ john.com: john.asm
 	$(CC) symlink.c -o ../run/keepass2john.exe
 	strip ../run/keepass2john.exe
 
+../run/keychain2john.exe: symlink.c
+	$(CC) symlink.c -o ../run/keychain2john.exe
+	strip ../run/keychain2john.exe
+
 ../run/pdf2john.exe: symlink.c
 	$(CC) symlink.c -o ../run/pdf2john.exe
 	strip ../run/pdf2john.exe
diff --git a/src/NS_fmt_plug.c b/src/NS_fmt_plug.c
index 87d813c..db86aca 100644
--- a/src/NS_fmt_plug.c
+++ b/src/NS_fmt_plug.c
@@ -113,8 +113,11 @@ static int NS_valid(char *ciphertext, struct fmt_main *self)
 
 static ARCH_WORD_32 *NS_std_get_binary(char *ciphertext)
 {
-	static unsigned long out_[BINARY_SIZE/sizeof(unsigned long)];
-	ARCH_WORD_32 *out = (ARCH_WORD_32*)out_;
+	static union {
+		unsigned long dummy;
+		ARCH_WORD_32 i[BINARY_SIZE/sizeof(ARCH_WORD_32)];
+	} _out;
+	ARCH_WORD_32 *out = _out.i;
 	char unscrambled[24];
 	int i;
         MD5_u32plus a, b, c;
diff --git a/src/alghmac.h b/src/alghmac.h
index 4ff87a2..6296887 100644
--- a/src/alghmac.h
+++ b/src/alghmac.h
@@ -34,13 +34,14 @@
 #ifndef _ALGHMAC_H_
 #define _ALGHMAC_H_
 
-#include "seccomon.h"
+#include <seccomon.h> /* NSS */
+/*
 #ifndef __APPLE__
 #include <malloc.h>
 #endif
 #include <memory.h>
 #include "alghmac.h"
-
+*/
 
 #define HMAC_PAD_SIZE 64
 
diff --git a/src/common-opencl.c b/src/common-opencl.c
index a09de07..f90b335 100644
--- a/src/common-opencl.c
+++ b/src/common-opencl.c
@@ -131,7 +131,7 @@ static void build_kernel(int dev_id)
 	if (build_code != CL_SUCCESS) {
 		// Give us much info about error and exit
 		fprintf(stderr, "Compilation log: %s\n", opencl_log);
-		fprintf(stderr, "Error building kernel. Returned build code: %d. DEVICE_INFO=%d\n", build_code, device_info[dev_id]); 
+		fprintf(stderr, "Error building kernel. Returned build code: %d. DEVICE_INFO=%d\n", build_code, device_info[dev_id]);
 		HANDLE_CLERROR (build_code, "clBuildProgram failed.");
 	}
 #ifdef REPORT_OPENCL_WARNINGS
diff --git a/src/cuda_mscash2_fmt.c b/src/cuda_mscash2_fmt.c
index 1b5179d..1ba0db2 100644
--- a/src/cuda_mscash2_fmt.c
+++ b/src/cuda_mscash2_fmt.c
@@ -168,6 +168,7 @@ static void *salt(char *ciphertext)
 	static mscash2_salt salt;
 	char *pos = ciphertext + strlen(mscash2_prefix);
 	int length = 0;
+	memset(&salt, 0, sizeof(salt));
 	salt.rounds = DEFAULT_ROUNDS;
 	sscanf(pos, "%d", &salt.rounds);
 	while (*pos++ != '#');
diff --git a/src/cuda_mscash_fmt.c b/src/cuda_mscash_fmt.c
index cb4e921..a846996 100644
--- a/src/cuda_mscash_fmt.c
+++ b/src/cuda_mscash_fmt.c
@@ -16,21 +16,21 @@
 
 #define FORMAT_LABEL		"mscash-cuda"
 #define FORMAT_NAME		"M$ Cache Hash MD4"
-#define ALGORITHM_NAME		"CUDA"
+#define ALGORITHM_NAME		"CUDA (inefficient, development use only)"
 #define MAX_CIPHERTEXT_LENGTH	(2 + 19*3 + 1 + 32)
-#define BENCHMARK_COMMENT	" len(pass)=8, len(salt)=13"
-#define BENCHMARK_LENGTH	-1
+#define BENCHMARK_COMMENT	""
+#define BENCHMARK_LENGTH	0
 
 static mscash_password *inbuffer;
 static mscash_hash *outbuffer;
 static mscash_salt currentsalt;
 
 static struct fmt_tests tests[] = {
+	{"M$test2#ab60bdb4493822b175486810ac2abe63", "test2"},
 	{"M$test1#64cd29e36a8431a2b111378564a10631", "test1"},
 	{"M$test1#64cd29e36a8431a2b111378564a10631", "test1"},
 	{"M$test1#64cd29e36a8431a2b111378564a10631", "test1"},
 	{"176a4c2bd45ac73687676c2f09045353", "", {"root"}},	// nullstring password
-	{"M$test2#ab60bdb4493822b175486810ac2abe63", "test2"},
 	{"M$test3#14dd041848e12fc48c0aa7a416a4a00c", "test3"},
 	{"M$test4#b945d24866af4b01a6d89b9d932a153c", "test4"},
 
@@ -122,6 +122,7 @@ static void *salt(char *ciphertext)
 	static mscash_salt salt;
 	char *pos = ciphertext + strlen(mscash_prefix);
 	int length = 0;
+	memset(&salt, 0, sizeof(salt));
 	while (*pos != '#') {
 		if (length == SALT_LENGTH)
 			return NULL;
diff --git a/src/cuda_phpass_fmt.c b/src/cuda_phpass_fmt.c
index 7b98e16..cd3ee38 100644
--- a/src/cuda_phpass_fmt.c
+++ b/src/cuda_phpass_fmt.c
@@ -16,7 +16,7 @@
 
 #define ALGORITHM_NAME		"CUDA"
 
-#define BENCHMARK_COMMENT	" ($P$9 lengths 1 to 15)"
+#define BENCHMARK_COMMENT	" ($P$9 lengths 0 to 15)"
 
 #define BENCHMARK_LENGTH	-1
 
@@ -36,33 +36,50 @@ static int any_cracked;
 extern void gpu_phpass(uint8_t *, phpass_salt *, phpass_crack *);
 
 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"},
-	{"$H$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$9saltstrikCftjZCE7EY2Kg/pjbl8S.", "abc"},
+	{"$P$900000000jPBDh/JWJIyrF0.DmP7kT.", "ala"},
+//	{"$P$9saltstriV/GXRIRi9UVeMLMph9BxF0", "abcd"},
+//	{"$P$900000000a94rg7R/nUK0icmALICKj1", "john"},
+//	{"$P$900000001ahWiA6cMRZxkgUxj4x/In0", "john"},
+//	{"$P$900000000a94rg7R/nUK0icmALICKj1", "john"},
+	{"$P$9sadli2.wzQIuzsR2nYVhUSlHNKgG/0", "john"},
 	{"$P$9saltstri3JPgLni16rBZtI03oeqT.0", "abcde"},
-	{"$P$9saltstri0D3A6JyITCuY72ZoXdejV.", "abcdef"},
+//	{"$P$9saltstri0D3A6JyITCuY72ZoXdejV.", "abcdef"},
+	{"$P$900000000zgzuX4Dc2091D8kak8RdR0", "h3ll00"},
 	{"$P$9saltstriXeNc.xV8N.K9cTs/XEn13.", "abcdefg"},
-	{"$P$9saltstrinwvfzVRP3u1gxG2gTLWqv.", "abcdefgh"},
-	{"$P$9saltstriSUQTD.yC2WigjF8RU0Q.Z.", "abcdefghi"},
-	{"$P$9saltstriWPpGLG.jwJkwGRwdKNEsg.", "abcdefghij"},
-
-	{"$P$9saltstrizjDEWUMXTlQHQ3/jhpR4C.", "abcdefghijk"},
+//	{"$P$9saltstrinwvfzVRP3u1gxG2gTLWqv.", "abcdefgh"},
+	{"$P$900000000m6YEJzWtTmNBBL4jypbHv1", "openwall"},
+	{"$H$9saltstriSUQTD.yC2WigjF8RU0Q.Z.", "abcdefghi"},
+//	{"$P$9saltstriWPpGLG.jwJkwGRwdKNEsg.", "abcdefghij"},
+//	{"$P$900000000qZTL5A0XQUX9hq0t8SoKE0", "1234567890"},
+	{"$P$900112200B9LMtPy2FSq910c1a6BrH0", "1234567890"},
+//	{"$P$9saltstrizjDEWUMXTlQHQ3/jhpR4C.", "abcdefghijk"},
+	{"$P$9RjH.g0cuFtd6TnI/A5MRR90TXPc43/", "password__1"},
 	{"$P$9saltstriGLUwnE6bl91BPJP6sxyka.", "abcdefghijkl"},
 	{"$P$9saltstriq7s97e2m7dXnTEx2mtPzx.", "abcdefghijklm"},
 	{"$P$9saltstriTWMzWKsEeiE7CKOVVU.rS0", "abcdefghijklmn"},
 	{"$P$9saltstriXt7EDPKtkyRVOqcqEW5UU.", "abcdefghijklmno"},
+#if 0
+	{"$H$9aaaaaSXBjgypwqm.JsMssPLiS8YQ00", "test1"},
+	{"$H$9PE8jEklgZhgLmZl5.HYJAzfGCQtzi1", "123456"},
+	{"$H$9pdx7dbOW3Nnt32sikrjAxYFjX8XoK1", "123456"},
+//	{"$P$912345678LIjjb6PhecupozNBmDndU0", "thisisalongertestPW"},
+	{"$H$9A5she.OeEiU583vYsRXZ5m2XIpI68/", "123456"},
+	{"$P$917UOZtDi6ksoFt.y2wUYvgUI6ZXIK/", "test1"},
+//	{"$P$91234567AQwVI09JXzrV1hEC6MSQ8I0", "thisisalongertest"},
+	{"$P$9234560A8hN6sXs5ir0NfozijdqT6f0", "test2"},
+	{"$P$9234560A86ySwM77n2VA/Ey35fwkfP0", "test3"},
+	{"$P$9234560A8RZBZDBzO5ygETHXeUZX5b1", "test4"},
+	{"$P$612345678si5M0DDyPpmRCmcltU/YW/", "JohnRipper"}, // 256
+	{"$P$6T4Krr44HLrUqGkL8Lu67lzZVbvHLC1", "test12345"}, // 256
+	{"$H$712345678WhEyvy1YWzT4647jzeOmo0", "JohnRipper"}, // 512 (phpBB w/older PHP version)
+	{"$P$8DkV/nqeaQNTdp4NvWjCkgN48AK69X.", "test12345"}, // 1024
+	{"$P$B12345678L6Lpt4BxNotVIMILOa9u81", "JohnRipper"}, // 8192 (WordPress)
+//	{"$P$91234567xogA.H64Lkk8Cx8vlWBVzH0", "thisisalongertst"},
+#endif
 	{NULL}
 };
 
diff --git a/src/cuda_pwsafe.h b/src/cuda_pwsafe.h
index 299e2cb..8728d2b 100644
--- a/src/cuda_pwsafe.h
+++ b/src/cuda_pwsafe.h
@@ -1,5 +1,5 @@
 /*
-* This software is Copyright (c) 2012 Lukas Odzioba <ukasz@openwall.net>
+* This software is Copyright (c) 2012 Lukas Odzioba <ukasz at openwall.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.
 */
diff --git a/src/cuda_rawsha256_fmt.c b/src/cuda_rawsha256_fmt.c
index 0240716..dab58dc 100644
--- a/src/cuda_rawsha256_fmt.c
+++ b/src/cuda_rawsha256_fmt.c
@@ -49,7 +49,7 @@ static struct fmt_tests sha224_tests[] = {
 	{NULL}
 };
 #endif
-#define ALGORITHM_NAME		"CUDA"
+#define ALGORITHM_NAME		"CUDA (inefficient, development use mostly)"
 
 extern void gpu_rawsha256(sha256_password *, SHA_HASH *, int);
 extern void gpu_rawsha224(sha256_password *, SHA_HASH *, int);
diff --git a/src/cuda_rawsha512_fmt.c b/src/cuda_rawsha512_fmt.c
index 920b058..21aef52 100644
--- a/src/cuda_rawsha512_fmt.c
+++ b/src/cuda_rawsha512_fmt.c
@@ -27,7 +27,7 @@
 
 #define FORMAT_LABEL			"raw-sha512-cuda"
 #define FORMAT_NAME			"Raw SHA-512"
-#define ALGORITHM_NAME			"CUDA"
+#define ALGORITHM_NAME			"CUDA (inefficient, development use mostly)"
 
 #define BENCHMARK_COMMENT		""
 #define BENCHMARK_LENGTH		-1
diff --git a/src/cuda_xsha512_fmt.c b/src/cuda_xsha512_fmt.c
index a90ec2e..0cdc37d 100644
--- a/src/cuda_xsha512_fmt.c
+++ b/src/cuda_xsha512_fmt.c
@@ -29,7 +29,7 @@
 
 #define FORMAT_LABEL			"xsha512-cuda"
 #define FORMAT_NAME			"Mac OS X 10.7+ salted SHA-512"
-#define ALGORITHM_NAME			"CUDA"
+#define ALGORITHM_NAME			"CUDA (efficient at \"many salts\" only)"
 
 #define BENCHMARK_COMMENT		""
 #define BENCHMARK_LENGTH		0
@@ -289,7 +289,7 @@ static void crypt_all(int count)
 
 static int cmp_all(void *binary, int count)
 {
-	int t1 = cuda_cmp_all(binary, count); 
+	int t1 = cuda_cmp_all(binary, count);
 	return t1;
 }
 
diff --git a/src/gladman_hmac.h b/src/gladman_hmac.h
index b6a6f8d..94fbdee 100644
--- a/src/gladman_hmac.h
+++ b/src/gladman_hmac.h
@@ -36,7 +36,7 @@
 #ifndef _HMAC_H
 #define _HMAC_H
 
-#include <memory.h>
+//#include <memory.h>
 
 #include "gladman_sha1.h"
 
diff --git a/src/gladman_pwd2key.c b/src/gladman_pwd2key.c
index 2eb2dc8..9756d29 100644
--- a/src/gladman_pwd2key.c
+++ b/src/gladman_pwd2key.c
@@ -36,7 +36,7 @@
  Compile: gcc -DTEST gladman_pwd2key.c gladman_hmac.c -lcrypto */
 
 #include <string.h>
-#include <memory.h>
+//#include <memory.h>
 #include "gladman_hmac.h"
 
 #if defined(__cplusplus)
diff --git a/src/john.c b/src/john.c
index 135fa21..abef357 100644
--- a/src/john.c
+++ b/src/john.c
@@ -190,8 +190,11 @@ extern int unshadow(int argc, char **argv);
 extern int unafs(int argc, char **argv);
 extern int undrop(int argc, char **argv);
 #ifndef _MSC_VER
+/* XXX: What's wrong with having these along with MSC? Perhaps this restriction
+ * was meant to apply to some of these only? Maybe SSH only? */
 extern int ssh2john(int argc, char **argv);
 extern int keepass2john(int argc, char **argv);
+extern int keychain2john(int argc, char **argv);
 extern int rar2john(int argc, char **argv);
 extern int racf2john(int argc, char **argv);
 extern int pwsafe2john(int argc, char **argv);
@@ -1336,6 +1339,11 @@ int main(int argc, char **argv)
 		return keepass2john(argc, argv);
 	}
 
+	if (!strcmp(name, "keychain2john")) {
+		CPU_detect_or_fallback(argv, 0);
+		return keychain2john(argc, argv);
+	}
+
 	if (!strcmp(name, "rar2john")) {
 		CPU_detect_or_fallback(argv, 0);
 		return rar2john(argc, argv);
diff --git a/src/keychain2john.c b/src/keychain2john.c
index ac3d5ed..74605fc 100644
--- a/src/keychain2john.c
+++ b/src/keychain2john.c
@@ -71,44 +71,54 @@ static void print_hex(unsigned char *str, int len)
 static void process_file(const char *filename)
 {
 	FILE *fp;
+	unsigned char buf[4];
+	long pos, cipheroff;
+	unsigned char salt[SALTLEN];
+	unsigned char iv[IVLEN];
+	unsigned char ct[CTLEN];
+	size_t nread;
+	int err;
 
 	if (!(fp = fopen(filename, "rb"))) {
 		fprintf(stderr, "! %s: %s\n", filename, strerror(errno));
 		return;
 	}
-	fseek(fp, -4, SEEK_END);
-	unsigned char buf[4];
+
+	err = fseek(fp, -4, SEEK_END);
 
 	while(1) {
-		fseek(fp, -8, SEEK_CUR);
-		if(fread(buf, 4, 1, fp) == 0) {
-			fprintf(stderr, "%s : Couldn't find db key. Is a keychain file?\n", filename);
+		err |= fseek(fp, -8, SEEK_CUR);
+		if (err || fread(buf, 4, 1, fp) != 1) {
+			fprintf(stderr, "%s: Couldn't find db key. Is a keychain file?\n", filename);
 			exit(1);
 		}
-		if(!memcmp(buf, magic, 4))
+		if (!memcmp(buf, magic, 4))
 			break;
 	}
 
-	long pos = ftell(fp) - 4;
+	pos = ftell(fp) - 4;
 
 	// ciphertext offset
-	fseek(fp, pos + 8, SEEK_SET);
-	long cipheroff = fget32(fp);
+	err = fseek(fp, pos + 8, SEEK_SET);
+	cipheroff = fget32(fp);
+	err |= ferror(fp);
 
 	// salt
-	fseek(fp, pos + 44, SEEK_SET);
-	unsigned char salt[SALTLEN];
-	fread(salt, SALTLEN, 1, fp);
+	err |= fseek(fp, pos + 44, SEEK_SET);
+	nread = fread(salt, SALTLEN, 1, fp);
 
 	// IV
-	fseek(fp, pos + 64, SEEK_SET);
-	unsigned char iv[IVLEN];
-	fread(iv, IVLEN, 1, fp);
+	err |= fseek(fp, pos + 64, SEEK_SET);
+	nread += fread(iv, IVLEN, 1, fp);
 
 	// ciphertext
-	fseek(fp, pos + cipheroff, SEEK_SET);
-	unsigned char ct[CTLEN];
-	fread(ct, CTLEN, 1, fp);
+	err |= fseek(fp, pos + cipheroff, SEEK_SET);
+	nread += fread(ct, CTLEN, 1, fp);
+
+	if (err || ferror(fp) || nread != 3) {
+		fprintf(stderr, "%s: Couldn't parse the file. Is a keychain file?\n", filename);
+		exit(1);
+	}
 
 	// output
 	printf("%s:$keychain$*", filename);
@@ -122,7 +132,7 @@ static void process_file(const char *filename)
 	fclose(fp);
 }
 
-int main(int argc, char **argv)
+int keychain2john(int argc, char **argv)
 {
 	int i;
 
diff --git a/src/lowpbe.c b/src/lowpbe.c
index 4791405..66103e5 100644
--- a/src/lowpbe.c
+++ b/src/lowpbe.c
@@ -35,7 +35,7 @@
 #include <plarena.h>
 #include "lowpbe.h"
 #include <openssl/sha.h>
-#include <memory.h>
+//#include <memory.h>
 #include <stdio.h>
 #include <hasht.h>
 #include <secoidt.h>
diff --git a/src/mozilla_des.h b/src/mozilla_des.h
index 6a3207d..b71ab19 100644
--- a/src/mozilla_des.h
+++ b/src/mozilla_des.h
@@ -37,9 +37,9 @@
 #define _DES_H_ 1
 
 //#include "blapit.h"
-#include "seccomon.h"
-#include "prtypes.h"
-#include <memory.h>
+#include <seccomon.h>
+#include <prtypes.h>
+//#include <memory.h>
 #include "KeyDBCracker.h"
 
 
diff --git a/src/mozilla_fmt.c b/src/mozilla_fmt.c
index 5990c13..6c48398 100644
--- a/src/mozilla_fmt.c
+++ b/src/mozilla_fmt.c
@@ -47,7 +47,7 @@
 #define FORMAT_NAME		"Mozilla SHA-1 3DES"
 #define ALGORITHM_NAME		"32/" ARCH_BITS_STR
 #define BENCHMARK_COMMENT	""
-#define BENCHMARK_LENGTH	0
+#define BENCHMARK_LENGTH	-1
 #define PLAINTEXT_LENGTH	16
 #define BINARY_SIZE		16
 #define SALT_SIZE		sizeof(*salt_struct)
diff --git a/src/mscash1_fmt_plug.c b/src/mscash1_fmt_plug.c
index 26625b7..104b6ae 100644
--- a/src/mscash1_fmt_plug.c
+++ b/src/mscash1_fmt_plug.c
@@ -38,11 +38,11 @@
 
 /* Note: some tests will be replaced in init() if running UTF-8 */
 static struct fmt_tests tests[] = {
+	{"M$test2#ab60bdb4493822b175486810ac2abe63", "test2" },
 	{"M$test1#64cd29e36a8431a2b111378564a10631", "test1" },
 	{"M$test1#64cd29e36a8431a2b111378564a10631", "test1" },
 	{"M$test1#64cd29e36a8431a2b111378564a10631", "test1" },
 	{"176a4c2bd45ac73687676c2f09045353", "", {"root"} }, // nullstring password
-	{"M$test2#ab60bdb4493822b175486810ac2abe63", "test2" },
 	{"M$test3#14dd041848e12fc48c0aa7a416a4a00c", "test3" },
 	{"M$test4#b945d24866af4b01a6d89b9d932a153c", "test4" },
 
diff --git a/src/mscash2_fmt_plug.c b/src/mscash2_fmt_plug.c
index caed760..937cee4 100644
--- a/src/mscash2_fmt_plug.c
+++ b/src/mscash2_fmt_plug.c
@@ -327,8 +327,11 @@ static void *get_salt(char *_ciphertext)
 
 static void *get_binary(char *ciphertext)
 {
-	static unsigned long out_[BINARY_SIZE/sizeof(unsigned long)];
-	unsigned int *out = (unsigned int*)out_;
+	static union {
+		unsigned long dummy;
+		unsigned int i[BINARY_SIZE/sizeof(unsigned int)];
+	} _out;
+	unsigned int *out = _out.i;
 	unsigned int i = 0;
 	unsigned int temp;
 
diff --git a/src/nt2_fmt_plug.c b/src/nt2_fmt_plug.c
index a9e99cc..b9afa1b 100644
--- a/src/nt2_fmt_plug.c
+++ b/src/nt2_fmt_plug.c
@@ -228,8 +228,11 @@ static char *prepare(char *split_fields[10], struct fmt_main *self)
 
 static void *binary(char *ciphertext)
 {
-	static unsigned long out_[DIGEST_SIZE/sizeof(unsigned long)];
-	unsigned int *out = (unsigned int*)out_;
+	static union {
+		unsigned long dummy;
+		unsigned int i[DIGEST_SIZE/sizeof(unsigned int)];
+	} _out;
+	unsigned int *out = _out.i;
 	unsigned int i;
 	unsigned int temp;
 
diff --git a/src/opencl/cryptmd5_kernel.cl b/src/opencl/cryptmd5_kernel.cl
index 43bbdb4..1ae39d3 100644
--- a/src/opencl/cryptmd5_kernel.cl
+++ b/src/opencl/cryptmd5_kernel.cl
@@ -1,5 +1,5 @@
 /*
-* This software is Copyright (c) 2011,2012 Lukas Odzioba <ukasz@openwall.com>  
+* This software is Copyright (c) 2011,2012 Lukas Odzioba <ukasz at openwall.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.
 */
diff --git a/src/opencl/cryptsha512_kernel_AMD.cl b/src/opencl/cryptsha512_kernel_AMD.cl
index b6176e4..23ccee9 100644
--- a/src/opencl/cryptsha512_kernel_AMD.cl
+++ b/src/opencl/cryptsha512_kernel_AMD.cl
@@ -306,7 +306,7 @@ inline void sha512crypt(__local  working_memory * fast_tmp_memory,
 __kernel
 // __attribute__((vec_type_hint(ulong2)))		Not recognized.
 // __attribute__((reqd_work_group_size(32, 1, 1)))	No gain.
-inline void kernel_crypt(__constant sha512_salt     * salt,
+void kernel_crypt(__constant sha512_salt     * salt,
                   __global   sha512_password * keys_buffer,
                   __global   sha512_hash     * out_buffer,
                   __local    sha512_salt     * salt_data,
@@ -359,4 +359,4 @@ inline void kernel_crypt(__constant sha512_salt     * salt,
 *   No gain, but the final performance was almost the same of this version.
 * - Tried to break this program into 2 kernels (prepare and crypt). Prepare do the job done
 *   outside the for loop. The gain was only 1%.
-***/
\ No newline at end of file
+***/
diff --git a/src/opencl/cryptsha512_kernel_CPU.cl b/src/opencl/cryptsha512_kernel_CPU.cl
index 6e40cc3..e381718 100644
--- a/src/opencl/cryptsha512_kernel_CPU.cl
+++ b/src/opencl/cryptsha512_kernel_CPU.cl
@@ -336,7 +336,7 @@ inline void sha512crypt(__global   sha512_salt     * salt_data,
 #undef pass
 
 __kernel
-inline void kernel_crypt(__global   sha512_salt     * informed_salt,
+void kernel_crypt(__global   sha512_salt     * informed_salt,
                   __global   sha512_password * pass_data,
                   __global   sha512_hash     * out_buffer) {
 
@@ -345,4 +345,4 @@ inline void kernel_crypt(__global   sha512_salt     * informed_salt,
 
     //Do the job
     sha512crypt(informed_salt,  &pass_data[gid], &out_buffer[gid]);
-}
\ No newline at end of file
+}
diff --git a/src/opencl/cryptsha512_kernel_DEFAULT.cl b/src/opencl/cryptsha512_kernel_DEFAULT.cl
index 2fe57f0..9d83d7e 100644
--- a/src/opencl/cryptsha512_kernel_DEFAULT.cl
+++ b/src/opencl/cryptsha512_kernel_DEFAULT.cl
@@ -309,7 +309,7 @@ inline void sha512_crypt(__global   sha512_salt     * salt_data,
 #undef ctx
 
 __kernel
-inline void kernel_crypt(__global   sha512_salt     * salt,
+void kernel_crypt(__global   sha512_salt     * salt,
                   __global   sha512_password * pass_data,
                   __global   sha512_hash     * out_buffer) {
 
@@ -328,4 +328,4 @@ inline void kernel_crypt(__global   sha512_salt     * salt,
     #pragma unroll
     for (int i = 0; i < 8; i++)
         out_buffer[gid].v[i] = buffers.alt_result[i].mem_64[0];
-}
\ No newline at end of file
+}
diff --git a/src/opencl/cryptsha512_kernel_NVIDIA.cl b/src/opencl/cryptsha512_kernel_NVIDIA.cl
index 9b53166..7486acc 100644
--- a/src/opencl/cryptsha512_kernel_NVIDIA.cl
+++ b/src/opencl/cryptsha512_kernel_NVIDIA.cl
@@ -339,7 +339,7 @@ inline void sha512crypt(__local  sha512_password * pass_data,
 #undef pass
 
 __kernel
-inline void kernel_crypt(__constant sha512_salt     * salt,
+void kernel_crypt(__constant sha512_salt     * salt,
                   __global   sha512_password * keys_buffer,
                   __global   sha512_hash     * out_buffer,
                   __local    sha512_salt     * salt_data,
@@ -354,4 +354,4 @@ inline void kernel_crypt(__constant sha512_salt     * salt,
 
     //Do the job
     sha512crypt(&tmp_memory[lid], salt_data, &out_buffer[gid]);
-}
\ No newline at end of file
+}
diff --git a/src/opencl/pbkdf2_kernel.cl b/src/opencl/pbkdf2_kernel.cl
index 4a374ac..ffa13ee 100644
--- a/src/opencl/pbkdf2_kernel.cl
+++ b/src/opencl/pbkdf2_kernel.cl
@@ -609,7 +609,7 @@ inline void hmac_sha1_iter(__private uint *istate, __private uint *ostate, __pri
 }
 
 __kernel 
-inline void PBKDF2 ( const __global unsigned int *pass_global, 
+void PBKDF2 ( const __global unsigned int *pass_global, 
               const __global unsigned int *salt, 
               int usrlen,  
               uint num_keys,
diff --git a/src/opencl/phpass_kernel.cl b/src/opencl/phpass_kernel.cl
index d446aaf..b5e33a2 100644
--- a/src/opencl/phpass_kernel.cl
+++ b/src/opencl/phpass_kernel.cl
@@ -1,295 +1,400 @@
 /*
-* 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.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
+
 typedef struct {
-	unsigned char v[PLAINTEXT_LENGTH];
-	unsigned char length;
+        uchar v[PLAINTEXT_LENGTH];
+        uchar length;
 } phpass_password;
 
 typedef struct {
-	unsigned int v[4];
+        uint v[4];
 } phpass_hash;
 
-#define ROTATE_LEFT(x, s)	rotate(x,(unsigned int)s)
 
-#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 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 F(x, y, z) bitselect((z), (y), (x))
+#define G(x, y, z) bitselect((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) += F ((b), (c), (d)) + (x) + (ac); \
+   (a) = ROTATE_LEFT ((a), (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) += G ((b), (c), (d)) + (x) + (ac); \
+   (a) = ROTATE_LEFT ((a), (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) += H ((b), (c), (d)) + (x) + (ac); \
+   (a) = ROTATE_LEFT ((a), (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); \
+  {(a) += I ((b), (c), (d)) + (x) + (ac); \
+   (a) = ROTATE_LEFT ((a), (s)) + (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
-
-inline void md5(char len,__private uint32_t * internal_ret,__private uint32_t * x)
+#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
+
+
+
+inline void md5(uint8 len,__private uint8 * internal_ret,__private uint8 * x)
 {
-	x[len / 4] |= (((uint32_t) 0x80) << ((len & 0x3) << 3));
-	uint32_t x14 = len << 3;
-
-	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;
+        uint8 x14 = len << 3;
+
+        uint8 a;
+        uint8 b = (uint8)0xefcdab89;
+        uint8 c = (uint8)0x98badcfe;
+        uint8 d = (uint8)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;
+        FF(a, b, c, d, x[4], S11,(uint8) 0xf57c0faf);
+        FF(d, a, b, c, x[5], S12,(uint8) 0x4787c62a);
+        FF(c, d, a, b, x[6], S13,(uint8) 0xa8304613);
+        FF(b, c, d, a, x[7], S14,(uint8) 0xfd469501);
+        FF(a, b, c, d, 0, S11,(uint8) 0x698098d8);
+        FF(d, a, b, c, 0, S12,(uint8) 0x8b44f7af);
+        FF(c, d, a, b, 0, S13,(uint8) 0xffff5bb1);
+        FF(b, c, d, a, 0, S14,(uint8) 0x895cd7be);
+        FF(a, b, c, d, 0, S11,(uint8) 0x6b901122);
+        FF(d, a, b, c, 0, S12,(uint8) 0xfd987193);
+        FF(c, d, a, b, x14, S13,(uint8) 0xa679438e);
+        FF(b, c, d, a, 0, S14,(uint8) 0x49b40821);
+
+        GG(a, b, c, d, x[1], S21,(uint8) 0xf61e2562);
+        GG(d, a, b, c, x[6], S22,(uint8) 0xc040b340);
+        GG(c, d, a, b, 0, S23,(uint8) 0x265e5a51);
+        GG(b, c, d, a, x[0], S24,(uint8) 0xe9b6c7aa);
+        GG(a, b, c, d, x[5], S21,(uint8) 0xd62f105d);
+        GG(d, a, b, c, 0, S22,(uint8) 0x2441453);
+        GG(c, d, a, b, 0, S23,(uint8) 0xd8a1e681);
+        GG(b, c, d, a, x[4], S24,(uint8) 0xe7d3fbc8);
+        GG(a, b, c, d, 0, S21,(uint8) 0x21e1cde6);
+        GG(d, a, b, c, x14, S22,(uint8) 0xc33707d6);
+        GG(c, d, a, b, x[3], S23,(uint8) 0xf4d50d87);
+        GG(b, c, d, a, 0, S24,(uint8) 0x455a14ed);
+        GG(a, b, c, d, 0, S21,(uint8) 0xa9e3e905);
+        GG(d, a, b, c, x[2], S22,(uint8) 0xfcefa3f8);
+        GG(c, d, a, b, x[7], S23,(uint8) 0x676f02d9);
+        GG(b, c, d, a, 0, S24,(uint8) 0x8d2a4c8a);
+
+        HH(a, b, c, d, x[5], S31,(uint8) 0xfffa3942);
+        HH(d, a, b, c, 0, S32,(uint8) 0x8771f681);
+        HH(c, d, a, b, 0, S33,(uint8) 0x6d9d6122);
+        HH(b, c, d, a, x14, S34,(uint8) 0xfde5380c);
+        HH(a, b, c, d, x[1], S31,(uint8) 0xa4beea44);
+        HH(d, a, b, c, x[4], S32,(uint8) 0x4bdecfa9);
+        HH(c, d, a, b, x[7], S33,(uint8) 0xf6bb4b60);
+        HH(b, c, d, a, 0, S34,(uint8) 0xbebfbc70);
+        HH(a, b, c, d, 0, S31,(uint8) 0x289b7ec6);
+        HH(d, a, b, c, x[0], S32,(uint8) 0xeaa127fa);
+        HH(c, d, a, b, x[3], S33,(uint8) 0xd4ef3085);
+        HH(b, c, d, a, x[6], S34,(uint8) 0x4881d05);
+        HH(a, b, c, d, 0, S31,(uint8) 0xd9d4d039);
+        HH(d, a, b, c, 0, S32,(uint8) 0xe6db99e5);
+        HH(c, d, a, b, 0, S33,(uint8) 0x1fa27cf8);
+        HH(b, c, d, a, x[2], S34,(uint8) 0xc4ac5665);
+
+        II(a, b, c, d, x[0], S41,(uint8) 0xf4292244);
+        II(d, a, b, c, x[7], S42,(uint8) 0x432aff97);
+        II(c, d, a, b, x14, S43,(uint8) 0xab9423a7);
+        II(b, c, d, a, x[5], S44,(uint8) 0xfc93a039);
+        II(a, b, c, d, 0, S41,(uint8) 0x655b59c3);
+        II(d, a, b, c, x[3], S42,(uint8) 0x8f0ccc92);
+        II(c, d, a, b, 0, S43,(uint8) 0xffeff47d);
+        II(b, c, d, a, x[1], S44,(uint8) 0x85845dd1);
+        II(a, b, c, d, 0, S41,(uint8) 0x6fa87e4f);
+        II(d, a, b, c, 0, S42,(uint8) 0xfe2ce6e0);
+        II(c, d, a, b, x[6], S43,(uint8) 0xa3014314);
+        II(b, c, d, a, 0, S44,(uint8) 0x4e0811a1);
+        II(a, b, c, d, x[4],S41,(uint8) 0xf7537e82);
+        II(d, a, b, c, 0, S42,(uint8) 0xbd3af235);
+        II(c, d, a, b, x[2], S43,(uint8) 0x2ad7d2bb);
+        II(b, c, d, a, 0, S44,(uint8) 0xeb86d391);
+
+        internal_ret[0] = a +(uint8) 0x67452301;
+        internal_ret[1] = b +(uint8) 0xefcdab89;
+        internal_ret[2] = c + (uint8)0x98badcfe;
+        internal_ret[3] = d +(uint8) 0x10325476;
 }
 
-inline void clear_ctx(__private uint32_t * x)
+inline void clear_ctx(__private uint8 * x)
 {
-	int i;
-	for (i = 0; i < 8; i++)
-		*x++ = 0;
+        uint8 zero = (uint8) (0, 0,0,0,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;
 }
 
+
 __kernel void phpass
     (   __global    const   phpass_password*    data
-    ,   __global            phpass_hash*    	data_out
-    ,   __global    const   char* 		setting
+    ,   __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;
-
-	uint32_t idx = get_global_id(0);
-
-	__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];
-	}
-
-	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;
-
-		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;
-
-	} 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
+        uint8 x[8],length;
+        uint sx[8],i,idx = get_global_id(0);
+        uint count = 1 << setting[SALT_SIZE+3];
+
+        clear_ctx(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;
+
+__private uchar *buff = (uchar *) sx;
+
+        #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);
+                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(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<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));
+
+
+
+#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; \
+ }
+
+uint8 a, b, c, d, x0, x1, x2, x3, x4, x5, x6, x7;
+uint8 x14 = len << 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 = (uint8)0xefcdab89;
+        c = (uint8)0x98badcfe;
+        d = (uint8)0x10325476;
+
+        a = AC1 + x0;
+        a = ROTATE_LEFT(a, S11);
+        a += b;                 /* 1 */
+        d = (c ^ (a & MASK1)) + x1 + AC2pCd;
+        d = ROTATE_LEFT(d, S12);
+        d += a;                 /* 2 */
+        c = F(d, a, b) + x2 + AC3pCc;
+        c = ROTATE_LEFT(c, S13);
+        c += d;                 /* 3 */
+        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);
+
+        x0 = a +(uint8) 0x67452301;
+        x1 = b +(uint8) 0xefcdab89;
+        x2 = c +(uint8)0x98badcfe;
+        x3 = d +(uint8) 0x10325476;
+        }while(--count);
+
+
+#define K3(q)\
+        data_out[idx*8+q].v[0]=x0.s##q;\
+        data_out[idx*8+q].v[1]=x1.s##q;\
+        data_out[idx*8+q].v[2]=x2.s##q;\
+        data_out[idx*8+q].v[3]=x3.s##q;
+
+        K3(0)
+        K3(1)
+        K3(2)
+        K3(3)
+        K3(4)
+        K3(5)
+        K3(6)
+        K3(7)
+
+#undef K3
+}
diff --git a/src/opencl/pwsafe_kernel.cl b/src/opencl/pwsafe_kernel.cl
index 47714ff..492bd2f 100644
--- a/src/opencl/pwsafe_kernel.cl
+++ b/src/opencl/pwsafe_kernel.cl
@@ -1,5 +1,5 @@
 /*
-* This software is Copyright (c) 2012 Lukas Odzioba <ukasz@openwall.net> 
+* This software is Copyright (c) 2012 Lukas Odzioba <ukasz at openwall.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.
 */
diff --git a/src/opencl/wpapsk_kernel.cl b/src/opencl/wpapsk_kernel.cl
index c06f5e0..fd5582a 100644
--- a/src/opencl/wpapsk_kernel.cl
+++ b/src/opencl/wpapsk_kernel.cl
@@ -1,5 +1,5 @@
 /*
-* This software is Copyright (c) 2012 Lukas Odzioba <ukasz@openwall.net> 
+* This software is Copyright (c) 2012 Lukas Odzioba <ukasz at openwall.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.
 */
diff --git a/src/opencl_cryptmd5_fmt.c b/src/opencl_cryptmd5_fmt.c
index d6e8524..ccc68dd 100644
--- a/src/opencl_cryptmd5_fmt.c
+++ b/src/opencl_cryptmd5_fmt.c
@@ -26,7 +26,7 @@
 #define FORMAT_NAME		"md5crypt"
 #define KERNEL_NAME		"cryptmd5"
 
-#define ALGORITHM_NAME		"OpenCL"
+#define ALGORITHM_NAME		"OpenCL (inefficient, development use mostly)"
 
 #define BENCHMARK_COMMENT	""
 #define BENCHMARK_LENGTH	-1
diff --git a/src/opencl_cryptsha512_fmt.c b/src/opencl_cryptsha512_fmt.c
index 3e775d7..f126ed0 100644
--- a/src/opencl_cryptsha512_fmt.c
+++ b/src/opencl_cryptsha512_fmt.c
@@ -297,7 +297,7 @@ static void find_best_workgroup(struct fmt_main *self) {
 /*  For a while reverted usage of common find_best_workgroup.
  */
     size_t max_group_size;
-        
+
     max_group_size = get_task_max_work_group_size();
     fprintf(stderr, "Max local work size %d, ", (int) max_group_size);
 
@@ -460,7 +460,6 @@ static void init(struct fmt_main *self) {
         task = "$JOHN/cryptsha512_kernel_CPU.cl";
 
     else {
-        fprintf(stderr, "Building the kernel, this could take a while\n");
         task = "$JOHN/cryptsha512_kernel_DEFAULT.cl";
 
         if (! no_byte_addressable(source_in_use)) {
@@ -471,6 +470,8 @@ static void init(struct fmt_main *self) {
                 task = "$JOHN/cryptsha512_kernel_AMD.cl";
         }
     }
+    fprintf(stderr, "Building the kernel (%s), this could take a while\n",
+	task);
     fflush(stdout);
     opencl_build_kernel(task, ocl_gpu_id);
     fflush(stdout);
diff --git a/src/opencl_mscash2_fmt.c b/src/opencl_mscash2_fmt.c
index 144981c..64eba71 100644
--- a/src/opencl_mscash2_fmt.c
+++ b/src/opencl_mscash2_fmt.c
@@ -60,13 +60,10 @@
     (((n) << 24) | (((n) & 0xff00) << 8) | (((n) >> 8) & 0xff00) | ((n) >> 24))
 
 
-typedef struct
-{
-	unsigned char username[MAX_SALT_LENGTH+1];
-
+typedef struct {
 	unsigned int length;
-
-} 	ms_cash2_salt;
+	unsigned char username[MAX_SALT_LENGTH+1];
+} ms_cash2_salt;
 
 
 
@@ -384,14 +381,13 @@ static void *binary(char *ciphertext)
 static void *salt(char *ciphertext)
 {
 	static ms_cash2_salt salt;
-
 	unsigned int length;
+	char *pos;
 
-	char *pos ;
+	memset(&salt, 0, sizeof(salt));
 
-	length=0;
-
-	pos=ciphertext + strlen(MSCASH2_PREFIX);
+	length = 0;
+	pos = ciphertext + strlen(MSCASH2_PREFIX);
 
 	while (*pos != '#')
 	      {
@@ -402,7 +398,6 @@ static void *salt(char *ciphertext)
 	      }
 
 	salt.username[length] = 0;
-
 	salt.length=length;
 
 	return &salt;
diff --git a/src/opencl_mysqlsha1_fmt.c b/src/opencl_mysqlsha1_fmt.c
index afc07fc..cf6a9cb 100644
--- a/src/opencl_mysqlsha1_fmt.c
+++ b/src/opencl_mysqlsha1_fmt.c
@@ -20,7 +20,7 @@
 
 #define FORMAT_LABEL			"mysql-sha1-opencl"
 #define FORMAT_NAME			"MySQL 4.1 double-SHA-1"
-#define ALGORITHM_NAME			"OpenCL"
+#define ALGORITHM_NAME			"OpenCL (inefficient, development use only)"
 
 #define BENCHMARK_COMMENT		""
 #define BENCHMARK_LENGTH		0
diff --git a/src/opencl_nsldaps_fmt.c b/src/opencl_nsldaps_fmt.c
index 1a8ab9e..b9e6613 100644
--- a/src/opencl_nsldaps_fmt.c
+++ b/src/opencl_nsldaps_fmt.c
@@ -31,7 +31,7 @@
 
 #define FORMAT_LABEL			"ssha-opencl"
 #define FORMAT_NAME			"Netscape LDAP salted SHA-1"
-#define ALGORITHM_NAME			"OpenCL"
+#define ALGORITHM_NAME			"OpenCL (inefficient, development use mostly)"
 
 #define BENCHMARK_COMMENT		""
 #define BENCHMARK_LENGTH		0
diff --git a/src/opencl_nt_fmt.c b/src/opencl_nt_fmt.c
index fee8c38..3932f80 100644
--- a/src/opencl_nt_fmt.c
+++ b/src/opencl_nt_fmt.c
@@ -38,6 +38,7 @@
 
 #define FORMAT_LABEL			"nt-opencl"
 #define FORMAT_NAME			"NT MD4"
+#define ALGORITHM_NAME			"OpenCL (inefficient, development use only)"
 
 #define BENCHMARK_COMMENT		""
 #define BENCHMARK_LENGTH		-1
@@ -120,7 +121,6 @@ static int have_full_hashes;
 
 static int max_keys_per_crypt = NT_NUM_KEYS;
 
-#define ALGORITHM_NAME		"OpenCL"
 #define NT_CRYPT_FUN		nt_crypt_all_opencl
 
 static void release_all(void)
diff --git a/src/opencl_phpass_fmt.c b/src/opencl_phpass_fmt.c
index 9991dc8..4c7a8ba 100644
--- a/src/opencl_phpass_fmt.c
+++ b/src/opencl_phpass_fmt.c
@@ -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.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,27 @@
 
 #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 lengths 0 to 15)"
+#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 ACTUAL_SALT_SIZE        8
+#define SALT_SIZE               (ACTUAL_SALT_SIZE + 1) // 1 byte for iterations
 
+//#define KEYS_PER_CRYPT          1024*8*40
+#define KEYS_PER_CRYPT          (2048 * 96)
+#define MIN_KEYS_PER_CRYPT      KEYS_PER_CRYPT
+#define MAX_KEYS_PER_CRYPT      KEYS_PER_CRYPT
 
 //#define _PHPASS_DEBUG
 
@@ -44,10 +45,11 @@ 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 const char phpass_prefix[] = "$P$";
-static char currentsalt[SALT_SIZE + 1];
+static phpass_password *inbuffer;			/** plaintext ciphertexts **/
+static phpass_hash *outbuffer;			/** calculated hashes **/
+static const char phpassP_prefix[] = "$P$";
+static const char phpassH_prefix[] = "$H$";
+static char currentsalt[SALT_SIZE];
 
 extern void mem_init(unsigned char *, uint32_t *, char *, char *, int);
 extern void mem_clear(void);
@@ -57,38 +59,54 @@ extern void gpu_phpass(void);
 static cl_mem mem_in, mem_out, mem_setting;
 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 settingsize = sizeof(uint8_t) * ACTUAL_SALT_SIZE + 4;
+//static size_t global_work_size = KEYS_PER_CRYPT/8;
 
 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$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$9saltstriWPpGLG.jwJkwGRwdKNEsg.", "abcdefghij"},
-	   {"$P$9saltstrizjDEWUMXTlQHQ3/jhpR4C.", "abcdefghijk"},
-	   {"$P$9saltstriGLUwnE6bl91BPJP6sxyka.", "abcdefghijkl"},
-	   {"$P$9saltstriq7s97e2m7dXnTEx2mtPzx.", "abcdefghijklm"},
-	   {"$P$9saltstriTWMzWKsEeiE7CKOVVU.rS0", "abcdefghijklmn"},
-	   {"$P$9saltstriXt7EDPKtkyRVOqcqEW5UU.", "abcdefghijklmno"},
-	 */
+	{"$P$90000000000tbNYOc9TwXvLEI62rPt1", ""},
+	{"$P$9saltstriAcRMGl.91RgbAD6WSq64z.", "a"},
+	{"$P$9saltstriMljTzvdluiefEfDeGGQEl/", "ab"},
+//	{"$P$9saltstrikCftjZCE7EY2Kg/pjbl8S.", "abc"},
+	{"$P$900000000jPBDh/JWJIyrF0.DmP7kT.", "ala"},
+//	{"$P$9saltstriV/GXRIRi9UVeMLMph9BxF0", "abcd"},
+//	{"$P$900000000a94rg7R/nUK0icmALICKj1", "john"},
+//	{"$P$900000001ahWiA6cMRZxkgUxj4x/In0", "john"},
+//	{"$P$900000000a94rg7R/nUK0icmALICKj1", "john"},
+	{"$P$9sadli2.wzQIuzsR2nYVhUSlHNKgG/0", "john"},
+	{"$P$9saltstri3JPgLni16rBZtI03oeqT.0", "abcde"},
+//	{"$P$9saltstri0D3A6JyITCuY72ZoXdejV.", "abcdef"},
+	{"$P$900000000zgzuX4Dc2091D8kak8RdR0", "h3ll00"},
+	{"$P$9saltstriXeNc.xV8N.K9cTs/XEn13.", "abcdefg"},
+//	{"$P$9saltstrinwvfzVRP3u1gxG2gTLWqv.", "abcdefgh"},
+	{"$P$900000000m6YEJzWtTmNBBL4jypbHv1", "openwall"},
+	{"$H$9saltstriSUQTD.yC2WigjF8RU0Q.Z.", "abcdefghi"},
+//	{"$P$9saltstriWPpGLG.jwJkwGRwdKNEsg.", "abcdefghij"},
+//	{"$P$900000000qZTL5A0XQUX9hq0t8SoKE0", "1234567890"},
+	{"$P$900112200B9LMtPy2FSq910c1a6BrH0", "1234567890"},
+//	{"$P$9saltstrizjDEWUMXTlQHQ3/jhpR4C.", "abcdefghijk"},
+	{"$P$9RjH.g0cuFtd6TnI/A5MRR90TXPc43/", "password__1"},
+	{"$P$9saltstriGLUwnE6bl91BPJP6sxyka.", "abcdefghijkl"},
+	{"$P$9saltstriq7s97e2m7dXnTEx2mtPzx.", "abcdefghijklm"},
+	{"$P$9saltstriTWMzWKsEeiE7CKOVVU.rS0", "abcdefghijklmn"},
+	{"$P$9saltstriXt7EDPKtkyRVOqcqEW5UU.", "abcdefghijklmno"},
+#if 0
+	{"$H$9aaaaaSXBjgypwqm.JsMssPLiS8YQ00", "test1"},
+	{"$H$9PE8jEklgZhgLmZl5.HYJAzfGCQtzi1", "123456"},
+	{"$H$9pdx7dbOW3Nnt32sikrjAxYFjX8XoK1", "123456"},
+//	{"$P$912345678LIjjb6PhecupozNBmDndU0", "thisisalongertestPW"},
+	{"$H$9A5she.OeEiU583vYsRXZ5m2XIpI68/", "123456"},
+	{"$P$917UOZtDi6ksoFt.y2wUYvgUI6ZXIK/", "test1"},
+//	{"$P$91234567AQwVI09JXzrV1hEC6MSQ8I0", "thisisalongertest"},
+	{"$P$9234560A8hN6sXs5ir0NfozijdqT6f0", "test2"},
+	{"$P$9234560A86ySwM77n2VA/Ey35fwkfP0", "test3"},
+	{"$P$9234560A8RZBZDBzO5ygETHXeUZX5b1", "test4"},
+	{"$P$612345678si5M0DDyPpmRCmcltU/YW/", "JohnRipper"}, // 256
+	{"$P$6T4Krr44HLrUqGkL8Lu67lzZVbvHLC1", "test12345"}, // 256
+	{"$H$712345678WhEyvy1YWzT4647jzeOmo0", "JohnRipper"}, // 512 (phpBB w/older PHP version)
+	{"$P$8DkV/nqeaQNTdp4NvWjCkgN48AK69X.", "test12345"}, // 1024
+	{"$P$B12345678L6Lpt4BxNotVIMILOa9u81", "JohnRipper"}, // 8192 (WordPress)
+//	{"$P$91234567xogA.H64Lkk8Cx8vlWBVzH0", "thisisalongertst"},
+#endif
 	{NULL}
 };
 
@@ -106,9 +124,10 @@ static void release_all(void)
 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 +140,12 @@ 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;
-
+	global_work_size = KEYS_PER_CRYPT / 8;
 	atexit(release_all);
-	opencl_init("$JOHN/phpass_kernel.cl", ocl_gpu_id,platform_id);
+	opencl_init("$JOHN/phpass_kernel.cl", ocl_gpu_id, platform_id);
 
 	/// Alocate memory
 	inbuffer =
@@ -136,9 +153,9 @@ 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_hash *) calloc(MAX_KEYS_PER_CRYPT, sizeof(phpass_hash));
 	assert(inbuffer != NULL);
+
 	mem_in =
 	    clCreateBuffer(context[ocl_gpu_id], CL_MEM_READ_ONLY, insize, NULL,
 	    &cl_error);
@@ -162,21 +179,24 @@ static void init(struct fmt_main *self)
 	HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 2, sizeof(mem_setting),
 		&mem_setting), "Error while setting mem_setting");
 
-	opencl_find_best_workgroup(self);
+//      opencl_find_best_workgroup(pFmt);
+
+	local_work_size = 64;
 }
 
-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;
@@ -227,7 +247,7 @@ static void *binary(char *ciphertext)
 
 static void *salt(char *ciphertext)
 {
-	static unsigned char salt[SALT_SIZE + 1];
+	static unsigned char salt[SALT_SIZE];
 	memcpy(salt, &ciphertext[4], 8);
 	salt[8] = ciphertext[3];
 	return salt;
@@ -236,19 +256,19 @@ static void *salt(char *ciphertext)
 
 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 };
+	char setting[SALT_SIZE + 3] = { 0 };
 	strcpy(setting, currentsalt);
-	strcpy(setting + SALT_SIZE, phpass_prefix);
-	setting[SALT_SIZE + 3] = atoi64[ARCH_INDEX(currentsalt[8])];
+	strcpy(setting + ACTUAL_SALT_SIZE, phpassP_prefix);
+	setting[ACTUAL_SALT_SIZE + 3] = atoi64[ARCH_INDEX(currentsalt[8])];
 	/// Copy data to gpu
 	HANDLE_CLERROR(clEnqueueWriteBuffer(queue[ocl_gpu_id], mem_in, CL_FALSE, 0,
 		insize, inbuffer, 0, NULL, NULL), "Copy data to gpu");
@@ -258,8 +278,8 @@ static void crypt_all(int count)
 
 	/// Run kernel
 	HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[ocl_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[ocl_gpu_id]), "clFinish");
 
 	/// Read the result back
@@ -273,11 +293,11 @@ static void crypt_all(int count)
 static int binary_hash_0(void *binary)
 {
 #ifdef _PHPASS_DEBUG
-	fprintf(stderr, "binary_hash_0 ");
+	printf("binary_hash_0 ");
 	int i;
 	uint32_t *b = binary;
 	for (i = 0; i < 4; i++)
-		fprintf(stderr, "%08x ", b[i]);
+		printf("%08x ", b[i]);
 	puts("");
 #endif
 	return (((ARCH_WORD_32 *) binary)[0] & 0xf);
@@ -316,10 +336,10 @@ static int binary_hash_6(void *binary)
 static int get_hash_0(int index)
 {
 #ifdef _PHPASS_DEBUG
-	fprintf(stderr, "get_hash_0:   ");
+	printf("get_hash_0:   ");
 	int i;
 	for (i = 0; i < 4; i++)
-		fprintf(stderr, "%08x ", outbuffer[index].v[i]);
+		printf("%08x ", outbuffer[index].v[i]);
 	puts("");
 #endif
 	return outbuffer[index].v[0] & 0xf;
@@ -370,7 +390,7 @@ static int cmp_all(void *binary, int count)
 	}
 #ifdef _PHPASS_DEBUG
 	puts("cmp_all = 0");
-#endif	/* _PHPASS_DEBUG */
+#endif				/* _PHPASS_DEBUG */
 	return 0;
 }
 
@@ -405,7 +425,7 @@ struct fmt_main fmt_opencl_phpass = {
 		BENCHMARK_LENGTH,
 		PLAINTEXT_LENGTH,
 		BINARY_SIZE,
-		SALT_SIZE + 1,
+		SALT_SIZE,
 		MIN_KEYS_PER_CRYPT,
 		MAX_KEYS_PER_CRYPT,
 		FMT_CASE | FMT_8_BIT,
diff --git a/src/opencl_pwsafe_fmt.c b/src/opencl_pwsafe_fmt.c
index 4e5481a..f3deeeb 100644
--- a/src/opencl_pwsafe_fmt.c
+++ b/src/opencl_pwsafe_fmt.c
@@ -1,7 +1,7 @@
 /* Password Safe cracker patch for JtR. Hacked together during May of
  * 2012 by Dhiru Kholia <dhiru.kholia at gmail.com>.
  *
- * OpenCL port by Lukas Odzioba <ukasz@openwall.net>
+ * OpenCL port by Lukas Odzioba <ukasz at openwall.net>
  *
  * This software is Copyright (c) 2012, Dhiru Kholia <dhiru.kholia at gmail.com>,
  * and it is hereby released to the general public under the following terms:
diff --git a/src/opencl_rawmd4_fmt.c b/src/opencl_rawmd4_fmt.c
index 4042f47..b1795eb 100644
--- a/src/opencl_rawmd4_fmt.c
+++ b/src/opencl_rawmd4_fmt.c
@@ -19,7 +19,7 @@
 #define PLAINTEXT_LENGTH    31
 #define FORMAT_LABEL        "raw-md4-opencl"
 #define FORMAT_NAME         "Raw MD4"
-#define ALGORITHM_NAME      "OpenCL"
+#define ALGORITHM_NAME      "OpenCL (inefficient, development use only)"
 #define BENCHMARK_COMMENT   ""
 #define BENCHMARK_LENGTH    -1
 #define CIPHERTEXT_LENGTH   32
diff --git a/src/opencl_rawmd5_fmt.c b/src/opencl_rawmd5_fmt.c
index c5666e0..881432d 100644
--- a/src/opencl_rawmd5_fmt.c
+++ b/src/opencl_rawmd5_fmt.c
@@ -19,7 +19,7 @@
 #define PLAINTEXT_LENGTH    31
 #define FORMAT_LABEL        "raw-md5-opencl"
 #define FORMAT_NAME         "Raw MD5"
-#define ALGORITHM_NAME      "OpenCL"
+#define ALGORITHM_NAME      "OpenCL (inefficient, development use only)"
 #define BENCHMARK_COMMENT   ""
 #define BENCHMARK_LENGTH    -1
 #define CIPHERTEXT_LENGTH   32
diff --git a/src/opencl_rawsha1_fmt.c b/src/opencl_rawsha1_fmt.c
index b8b2819..198a3b6 100644
--- a/src/opencl_rawsha1_fmt.c
+++ b/src/opencl_rawsha1_fmt.c
@@ -19,8 +19,8 @@
 #include "common-opencl.h"
 
 #define FORMAT_LABEL			"raw-sha1-opencl"
-#define FORMAT_NAME			"Raw SHA-1 OpenCL"
-#define ALGORITHM_NAME			"OpenCL"
+#define FORMAT_NAME			"Raw SHA-1"
+#define ALGORITHM_NAME			"OpenCL (inefficient, development use only)"
 
 #define BENCHMARK_COMMENT		""
 #define BENCHMARK_LENGTH		-1
diff --git a/src/opencl_rawsha512_fmt.c b/src/opencl_rawsha512_fmt.c
index 917ab08..fd5ef3e 100644
--- a/src/opencl_rawsha512_fmt.c
+++ b/src/opencl_rawsha512_fmt.c
@@ -4,7 +4,7 @@
  * compares with ciphertext d80. For more details, refer to:
  * http://www.openwall.com/lists/john-dev/2012/04/11/13
  *
- * Copyright (c) 2012 myrice (interfacing to CUDA)
+ * Copyright (c) 2012 myrice
  *
  * Redistribution and use in source and binary forms, with or without
  * modification, are permitted.
@@ -28,7 +28,7 @@
 
 #define FORMAT_LABEL			"raw-sha512-opencl"
 #define FORMAT_NAME			"Raw SHA-512"
-#define ALGORITHM_NAME			"OpenCL"
+#define ALGORITHM_NAME			"OpenCL (inefficient, development use mostly)"
 
 #define BENCHMARK_COMMENT		""
 #define BENCHMARK_LENGTH		-1
@@ -94,7 +94,7 @@ static struct fmt_tests tests[] = {
 
 static sha512_key gkey[MAX_KEYS_PER_CRYPT];
 static sha512_hash ghash[MAX_KEYS_PER_CRYPT];
-uint8_t sha512_key_changed;
+static uint8_t sha512_key_changed;
 static uint8_t hash_copy_back;
 
 static uint64_t H[8] = {
@@ -448,5 +448,3 @@ struct fmt_main fmt_opencl_rawsha512 = {
 #warning Note: Mac OS X Lion format disabled - it needs OpenSSL 0.9.8 or above
 #endif
 #endif
-
-
diff --git a/src/opencl_wpapsk_fmt.c b/src/opencl_wpapsk_fmt.c
index 2d69e11..d795f68 100644
--- a/src/opencl_wpapsk_fmt.c
+++ b/src/opencl_wpapsk_fmt.c
@@ -1,5 +1,5 @@
 /*
-* This software is Copyright (c) 2012 Lukas Odzioba <ukasz@openwall.net>
+* This software is Copyright (c) 2012 Lukas Odzioba <ukasz at openwall.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.
 *
diff --git a/src/opencl_xsha512_fmt.c b/src/opencl_xsha512_fmt.c
index 206b24f..8df0eda 100644
--- a/src/opencl_xsha512_fmt.c
+++ b/src/opencl_xsha512_fmt.c
@@ -28,7 +28,7 @@
 
 #define FORMAT_LABEL			"xsha512-opencl"
 #define FORMAT_NAME			"Mac OS X 10.7+ salted SHA-512"
-#define ALGORITHM_NAME			"OpenCL"
+#define ALGORITHM_NAME			"OpenCL (efficient at \"many salts\" only)"
 
 #define BENCHMARK_COMMENT		""
 #define BENCHMARK_LENGTH		0
@@ -64,7 +64,7 @@
 #define FULL_BINARY_SIZE 64
 
 
-#define PLAINTEXT_LENGTH 20	
+#define PLAINTEXT_LENGTH 20
 #define CIPHERTEXT_LENGTH 136
 
 typedef struct {		// notice memory align problem
diff --git a/src/options.c b/src/options.c
index 5a24f89..4cb4f19 100644
--- a/src/options.c
+++ b/src/options.c
@@ -179,7 +179,7 @@ static struct opt_entry opt_list[] = {
 "                          For a full list of NAME use --list=encodings\n" \
 "--rules[=SECTION]         enable word mangling rules for wordlist modes\n" \
 "--incremental[=MODE]      \"incremental\" mode [using section MODE]\n" \
-"--markov[=options]        \"Markov\" mode (see doc/MARKOV)\n" \
+"--markov[=OPTIONS]        \"Markov\" mode (see doc/MARKOV)\n" \
 "--external=MODE           external mode or word filter\n" \
 "--stdout[=LENGTH]         just output candidate passwords [cut at LENGTH]\n" \
 "--restore[=NAME]          restore an interrupted session [called NAME]\n" \
diff --git a/src/rawSHA1_ng_fmt.c b/src/rawSHA1_ng_fmt.c
index 6cd42ae..5f89cda 100644
--- a/src/rawSHA1_ng_fmt.c
+++ b/src/rawSHA1_ng_fmt.c
@@ -51,7 +51,7 @@
 #define SHA1_DIGEST_WORDS        5
 #define SHA1_PARALLEL_HASH     512 // This must be a multiple of 4.
 
-#define __aligned __attribute__((aligned(16)))
+#define __aligned_16 __attribute__((aligned(16)))
 
 #ifndef __XOP__
 # define _mm_slli_epi32a(a, s)                                                                          \
@@ -142,12 +142,12 @@
 // M and N contain the first and last 128bits of a 512bit SHA-1 message block
 // respectively. The remaining 256bits are always zero, and so are not stored
 // here to avoid the load overhead.
-static uint32_t __aligned M[SHA1_PARALLEL_HASH][4];
-static uint32_t __aligned N[SHA1_PARALLEL_HASH][4];
+static uint32_t __aligned_16 M[SHA1_PARALLEL_HASH][4];
+static uint32_t __aligned_16 N[SHA1_PARALLEL_HASH][4];
 
 // MD contains the state of the SHA-1 A register at R75 for each of the input
 // messages.
-static uint32_t __aligned MD[SHA1_PARALLEL_HASH];
+static uint32_t __aligned_16 MD[SHA1_PARALLEL_HASH];
 
 static const char kFormatTag[] = "$dynamic_26$";
 
@@ -301,7 +301,7 @@ static void sha1_fmt_set_key(char *key, int index)
     // Create a lookup tables to find correct masks for each supported input
     // length. It would be nice if could use 128 bit shifts to produce these
     // dynamically, but they require an immediate operand.
-    static const __aligned uint32_t kTrailingBitTable[][4] = {
+    static const __aligned_16 uint32_t kTrailingBitTable[][4] = {
         { 0x00000080, 0x00000000, 0x00000000, 0x00000000 },
         { 0x00008000, 0x00000000, 0x00000000, 0x00000000 },
         { 0x00800000, 0x00000000, 0x00000000, 0x00000000 },
@@ -320,7 +320,7 @@ static void sha1_fmt_set_key(char *key, int index)
         { 0x00000000, 0x00000000, 0x00000000, 0x80000000 },
     };
 
-    static const __aligned uint32_t kUsedBytesTable[][4] = {
+    static const __aligned_16 uint32_t kUsedBytesTable[][4] = {
         { 0xFFFFFFFF, 0xFFFFFFFF, 0xFFFFFFFF, 0xFFFFFFFF },
         { 0xFFFFFF00, 0xFFFFFFFF, 0xFFFFFFFF, 0xFFFFFFFF },
         { 0xFFFF0000, 0xFFFFFFFF, 0xFFFFFFFF, 0xFFFFFFFF },
@@ -553,7 +553,7 @@ static inline int _mm_testz_epi32 (__m128i __X)
 # warning not using optimized sse4.1 compare because -msse4 was not specified
 static inline int _mm_testz_epi32 (__m128i __X)
 {
-    uint32_t __aligned words[4];
+    uint32_t __aligned_16 words[4];
     _mm_store_si128(words, __X);
     return !words[0] || !words[1] || !words[2] || !words[3];
 }
diff --git a/src/rawmd5u_fmt_plug.c b/src/rawmd5u_fmt_plug.c
index 5ecdb61..d45a3c0 100644
--- a/src/rawmd5u_fmt_plug.c
+++ b/src/rawmd5u_fmt_plug.c
@@ -165,8 +165,11 @@ static int valid(char *ciphertext, struct fmt_main *self)
 
 static void *binary(char *ciphertext)
 {
-	static unsigned long out_[BINARY_SIZE/sizeof(unsigned long)];
-	unsigned int *out = (unsigned int*)out_;
+	static union {
+		unsigned long dummy;
+		unsigned int i[BINARY_SIZE/sizeof(unsigned int)];
+	} _out;
+	unsigned int *out = _out.i;
 	unsigned int i;
 	unsigned int temp;
 
diff --git a/src/recovery.c b/src/recovery.c
index 4898731..baabc50 100644
--- a/src/recovery.c
+++ b/src/recovery.c
@@ -5,7 +5,11 @@
  * ...with changes in the jumbo patch, by JimF.
  */
 
+#ifndef __FreeBSD__
+/* On FreeBSD, defining this precludes the declaration of u_int, which
+ * FreeBSD's own <sys/file.h> needs. */
 #define _XOPEN_SOURCE 500 /* for fdopen(3), fileno(3), fsync(2), ftruncate(2) */
+#endif
 
 #include <stdio.h>
 #ifndef _MSC_VER
diff --git a/src/sip_fmt_plug.c b/src/sip_fmt_plug.c
index 3d33639..3958240 100644
--- a/src/sip_fmt_plug.c
+++ b/src/sip_fmt_plug.c
@@ -82,7 +82,7 @@ static int valid(char *ciphertext, struct fmt_main *self)
 }
 
 // NOTE, this still needs work. I am sure this will not eliminate (compact out)
-// duplicate salts. 
+// duplicate salts.
 static void *get_salt(char *ciphertext)
 {
 	static sip_salt salt;
diff --git a/src/stdint.h b/src/stdint.h
index 9b6ddc1..9ef75b3 100644
--- a/src/stdint.h
+++ b/src/stdint.h
@@ -1,5 +1,5 @@
 #ifndef _STDINT_H
-#ifdef __STDC_VERSION__
+#if defined(__STDC__) || defined(__STDC_VERSION__)
 #include <stdint.h>
 #else
 #define _STDINT_H 1
diff --git a/src/unrar.c b/src/unrar.c
index 2f103e3..378f04f 100644
--- a/src/unrar.c
+++ b/src/unrar.c
@@ -546,10 +546,10 @@ static int read_tables(const unsigned char **fd, unpack_data_t *unpack_data)
 		rar_dbgmsg("ERROR: read_tables check failed\n");
 		return 0;
 	}
-	rar_make_decode_tables(&table[0], (struct Decode *)&unpack_data->LD,NC);
-	rar_make_decode_tables(&table[NC], (struct Decode *)&unpack_data->DD,DC);
-	rar_make_decode_tables(&table[NC+DC], (struct Decode *)&unpack_data->LDD,LDC);
-	rar_make_decode_tables(&table[NC+DC+LDC], (struct Decode *)&unpack_data->RD,RC);
+	rar_make_decode_tables(&table[0], &unpack_data->LD.D,NC);
+	rar_make_decode_tables(&table[NC], &unpack_data->DD.D,DC);
+	rar_make_decode_tables(&table[NC+DC], &unpack_data->LDD.D,LDC);
+	rar_make_decode_tables(&table[NC+DC+LDC], &unpack_data->RD.D,RC);
 	memcpy(unpack_data->unp_old_table,table,sizeof(unpack_data->unp_old_table));
 
 
diff --git a/src/unrar.h b/src/unrar.h
index c604dfd..79817ee 100644
--- a/src/unrar.h
+++ b/src/unrar.h
@@ -161,11 +161,26 @@ typedef struct unpack_data_tag
 	int prev_low_dist;
 	int low_dist_rep_count;
 	unsigned char unp_old_table[HUFF_TABLE_SIZE];
-	struct LitDecode LD;
-	struct DistDecode DD;
-	struct LowDistDecode LDD;
-	struct RepDecode RD;
-	struct BitDecode BD;
+	union {
+		struct LitDecode LD;
+		struct Decode D;
+	} LD;
+	union {
+		struct DistDecode DD;
+		struct Decode D;
+	} DD;
+	union {
+		struct LowDistDecode LDD;
+		struct Decode D;
+	} LDD;
+	union {
+		struct RepDecode RD;
+		struct Decode D;
+	} RD;
+	union {
+		struct BitDecode BD;
+		struct Decode D;
+	} BD;
 	unsigned int old_dist[4];
 	unsigned int old_dist_ptr;
 	unsigned int last_dist;