diff -urpN magnum-jumbo/src/cuda/cryptmd5.cu magnum-jumbo_cudatsfix//src/cuda/cryptmd5.cu
--- magnum-jumbo/src/cuda/cryptmd5.cu	2012-06-24 21:19:23.800587954 +0000
+++ magnum-jumbo_cudatsfix//src/cuda/cryptmd5.cu	2012-06-24 21:28:39.133826803 +0000
@@ -1,5 +1,5 @@
 /*
-* This software is Copyright (c) 2011 Lukas Odzioba <lukas dot odzioba at gmail dot 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.
 */
@@ -131,8 +131,7 @@ __device__ void md5_digest(md5_ctx * ctx
 }
 
 
-__device__ void md5crypt(const char *gpass, size_t keysize, uint32_t * result,
-    uint32_t idx)
+__device__ void md5crypt(const char *gpass, size_t keysize, char *result)
 {
 
 	uint32_t i;
@@ -141,21 +140,21 @@ __device__ void md5crypt(const char *gpa
 
 	uint8_t ctx_buflen;
 	char *pass = spass[threadIdx.x];
-	memcpy(pass, gpass, 16);
+	memcpy(pass, gpass, 15);
 	uint8_t pass_len = keysize;
-	uint8_t salt_len = cuda_salt[0].saltlen;
+	uint8_t salt_len = cuda_salt[0].length;
 	char *salt = cuda_salt[0].salt;
 	md5_ctx ctx;
 	ctx_init(&ctx, &ctx_buflen);
-	
+
 	ctx_update(&ctx, pass, pass_len, &ctx_buflen);
 	ctx_update(&ctx, salt, salt_len, &ctx_buflen);
 	ctx_update(&ctx, pass, pass_len, &ctx_buflen);
 	md5_digest(&ctx, alt_result[threadIdx.x], &ctx_buflen);
-		
+
 	ctx_init(&ctx, &ctx_buflen);
 
-		
+
 	ctx_update(&ctx, pass, pass_len, &ctx_buflen);
 	if (cuda_salt[0].prefix == '1') {
 		ctx_update(&ctx, md5_salt_prefix_cu, 3, &ctx_buflen);
@@ -163,7 +162,7 @@ __device__ void md5crypt(const char *gpa
 		ctx_update(&ctx, apr1_salt_prefix_cu, 6, &ctx_buflen);
 
 	ctx_update(&ctx, salt, salt_len, &ctx_buflen);
-	
+
 
 	ctx_update(&ctx, (const char *) alt_result[threadIdx.x], pass_len,
 	    &ctx_buflen);
@@ -202,21 +201,22 @@ __device__ void md5crypt(const char *gpa
 		else
 			ctx_update(&ctx, pass, pass_len, &ctx_buflen);
 		md5_digest(&ctx, alt_result[threadIdx.x], &ctx_buflen);
-	}
-
-	result[address(0, idx)] = alt_result[threadIdx.x][0];
-	result[address(1, idx)] = alt_result[threadIdx.x][1];
-	result[address(2, idx)] = alt_result[threadIdx.x][2];
-	result[address(3, idx)] = alt_result[threadIdx.x][3];
+	}	
+	char cracked = 1;
+	cracked &= (alt_result[threadIdx.x][0] == cuda_salt[0].hash[0]);
+	cracked &= (alt_result[threadIdx.x][1] == cuda_salt[0].hash[1]);
+	cracked &= (alt_result[threadIdx.x][2] == cuda_salt[0].hash[2]);
+	cracked &= (alt_result[threadIdx.x][3] == cuda_salt[0].hash[3]);
+	*result = cracked;
 }
 
 
 __global__ void kernel_crypt_r(crypt_md5_password * inbuffer,
-    uint32_t * outbuffer)
+    crypt_md5_crack * outbuffer)
 {
 	uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
 	md5crypt((char *) inbuffer[idx].v, inbuffer[idx].length,
-	    outbuffer, idx);
+	    &outbuffer[idx].cracked);
 }
 
 __host__ void md5_crypt_gpu(crypt_md5_password * inbuffer,
@@ -225,18 +225,19 @@ __host__ void md5_crypt_gpu(crypt_md5_pa
 	HANDLE_ERROR(cudaMemcpyToSymbol(cuda_salt, host_salt,
 		sizeof(crypt_md5_salt)));
 	crypt_md5_password *cuda_inbuffer;
-	uint32_t *cuda_outbuffer;
+	crypt_md5_crack *cuda_outbuffer;
 
 	size_t insize = sizeof(crypt_md5_password) * KEYS_PER_CRYPT;
-	size_t outsize = sizeof(crypt_md5_hash) * KEYS_PER_CRYPT;
+	size_t outsize = sizeof(crypt_md5_crack) * KEYS_PER_CRYPT;
 
 	HANDLE_ERROR(cudaMalloc(&cuda_inbuffer, insize));
 	HANDLE_ERROR(cudaMalloc(&cuda_outbuffer, outsize));
+
+
 	HANDLE_ERROR(cudaMemcpy(cuda_inbuffer, inbuffer, insize,
 		cudaMemcpyHostToDevice));
 
-	kernel_crypt_r <<< BLOCKS, THREADS >>> (cuda_inbuffer,
-	    cuda_outbuffer);
+	kernel_crypt_r <<< BLOCKS, THREADS >>> (cuda_inbuffer, cuda_outbuffer);
 
 	HANDLE_ERROR(cudaMemcpy(outbuffer, cuda_outbuffer, outsize,
 		cudaMemcpyDeviceToHost));
diff -urpN magnum-jumbo/src/cuda/phpass.cu magnum-jumbo_cudatsfix//src/cuda/phpass.cu
--- magnum-jumbo/src/cuda/phpass.cu	2012-06-24 21:19:23.800587954 +0000
+++ magnum-jumbo_cudatsfix//src/cuda/phpass.cu	2012-06-20 02:46:23.000000000 +0000
@@ -1,5 +1,5 @@
 /*
-* This software is Copyright (c) 2011 Lukas Odzioba <lukas dot odzioba at gmail dot 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.
 */
@@ -10,56 +10,32 @@
 #include "../cuda_phpass.h"
 #include "cuda_common.cuh"
 
-const uint SETTING_SIZE = 12;
 const uint DATA_IN_SIZE = KEYS_PER_CRYPT * sizeof(phpass_password);
-const uint DATA_OUT_SIZE = KEYS_PER_CRYPT * sizeof(phpass_hash);
+const uint DATA_OUT_SIZE = KEYS_PER_CRYPT * sizeof(phpass_crack);
 
-__device__ unsigned char *cuda_data = NULL;
-__device__ uint32_t *cuda_data_out = NULL;
-__device__ char *cuda_setting = NULL;
+__device__ __constant__ phpass_salt cuda_salt[1];
 
-unsigned char *host_data = NULL;
-uint32_t *host_data_out = NULL;
+__global__ void kernel_phpass(unsigned char *, phpass_crack *);
 
-extern "C" void mem_init(unsigned char *, uint32_t *, char *, char *, int);
-extern "C" void mem_clear(void);
-extern "C" void gpu_phpass(void);
-
-__global__ void kernel_phpass(unsigned char *data, uint32_t * data_out,
-    char *, int);
-
-int cuda_count_log2;
-__host__ void gpu_phpass(void)
+extern "C" void gpu_phpass(uint8_t * host_data, phpass_salt * salt,
+    phpass_crack * host_data_out)
 {
-	dim3 dimGrid(BLOCKS);
-	dim3 dimBlock(THREADS);
-	kernel_phpass <<< dimGrid, dimBlock >>> (cuda_data, cuda_data_out,
-	    cuda_setting, cuda_count_log2);
-	HANDLE_ERROR(cudaThreadSynchronize());
-}
-
-__host__ void mem_init(unsigned char *p, uint32_t * h, char *setting,
-    char *itoa, int count_log2)
-{
-	cuda_count_log2 = count_log2;
-	host_data = p;
-	host_data_out = h;
-	HANDLE_ERROR(cudaMalloc(&cuda_setting, SETTING_SIZE));
+	uint8_t *cuda_data;
+	phpass_crack *cuda_data_out;
 	HANDLE_ERROR(cudaMalloc(&cuda_data, DATA_IN_SIZE));
 	HANDLE_ERROR(cudaMalloc(&cuda_data_out, DATA_OUT_SIZE));
+
 	HANDLE_ERROR(cudaMemcpy(cuda_data, host_data, DATA_IN_SIZE,
 		cudaMemcpyHostToDevice));
-	HANDLE_ERROR(cudaMemcpy(cuda_setting, setting, SETTING_SIZE,
-		cudaMemcpyHostToDevice));
-}
+	HANDLE_ERROR(cudaMemcpyToSymbol(cuda_salt, salt, SALT_SIZE));
 
-__host__ void mem_clear()
-{
+	kernel_phpass <<< BLOCKS, THREADS >>> (cuda_data, cuda_data_out);
+
+	HANDLE_ERROR(cudaThreadSynchronize());
 	HANDLE_ERROR(cudaMemcpy(host_data_out, cuda_data_out, DATA_OUT_SIZE,
 		cudaMemcpyDeviceToHost));
 	HANDLE_ERROR(cudaFree(cuda_data));
 	HANDLE_ERROR(cudaFree(cuda_data_out));
-	HANDLE_ERROR(cudaFree(cuda_setting));
 }
 
 __device__ void cuda_md5(char len, uint32_t * internal_ret, uint32_t * x)
@@ -162,8 +138,7 @@ __device__ void clear_ctx(uint32_t * x)
 		*x++ = 0;
 }
 
-__global__ void kernel_phpass(unsigned char *password, uint32_t * data_out,
-    char *setting, int count_log2)
+__global__ void kernel_phpass(unsigned char *password, phpass_crack * data_out)
 {
 	uint32_t x[8];
 	clear_ctx(x);
@@ -176,18 +151,17 @@ __global__ void kernel_phpass(unsigned c
 
 #pragma unroll 8
 	for (i = 0; i < 8; i++)
-		buff[i] = setting[i + 4];
+		buff[i] = cuda_salt[0].salt[i];
 
 	for (i = 8; i < 8 + length; i++) {
 		buff[i] = password[address(i - 8, idx)];
 	}
 
 	cuda_md5(8 + length, x, x);
-	count = 1 << count_log2;
+	count = cuda_salt[0].rounds;
 	for (i = 16; i < 16 + length; i++)
 		buff[i] = password[address(i - 16, idx)];
 
-
 	uint32_t a, b, c, d, x0, x1, x2, x3, x4, x5, x6, x7;
 	uint32_t len = 16 + length;
 	uint32_t x14 = len << 3;
@@ -290,9 +264,10 @@ __global__ void kernel_phpass(unsigned c
 
 	} while (--count);
 
-	data_out[address(0, idx)] = x0;
-	data_out[address(1, idx)] = x1;
-	data_out[address(2, idx)] = x2;
-	data_out[address(3, idx)] = x3;
-
+	char cracked = 1;
+	cracked &= (x0 == cuda_salt[0].hash[0]);
+	cracked &= (x1 == cuda_salt[0].hash[1]);
+	cracked &= (x2 == cuda_salt[0].hash[2]);
+	cracked &= (x3 == cuda_salt[0].hash[3]);
+	data_out[idx].cracked = cracked;
 }
diff -urpN magnum-jumbo/src/cuda_cryptmd5.h magnum-jumbo_cudatsfix//src/cuda_cryptmd5.h
--- magnum-jumbo/src/cuda_cryptmd5.h	2012-06-24 21:19:23.800587954 +0000
+++ magnum-jumbo_cudatsfix//src/cuda_cryptmd5.h	2012-06-20 03:09:18.000000000 +0000
@@ -1,14 +1,14 @@
 /*
-* This software is Copyright (c) 2011 Lukas Odzioba <lukas dot odzioba at gmail dot 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.
 */
-#ifndef _CRYPTMD5_H
-#define _CRYPTMD5_H
+#ifndef _CUDA_CRYPTMD5_H
+#define _CUDA_CRYPTMD5_H
 #include <assert.h>
-#include "common.h"
 #include <stdint.h>
 #include <stdbool.h>
+#include "common.h"
 
 #define uint32_t unsigned int
 #define uint8_t unsigned char
@@ -16,25 +16,26 @@
 #define MIN(a,b) ((a)<(b)?(a):(b))
 #define MAX(a,b) ((a)>(b)?(a):(b))
 
-#define BLOCKS (28)
-#define THREADS 256 
-#define KEYS_PER_CRYPT BLOCKS*THREADS
+#define BLOCKS			28*3
+#define THREADS 		256 
+#define KEYS_PER_CRYPT		BLOCKS*THREADS
 #define PLAINTEXT_LENGTH	15
 
 typedef struct {
-	unsigned char saltlen;
+	uint32_t hash[4];	//hash that we are looking for
+	uint8_t length;   //salt length
 	char salt[8];
 	char prefix;		// 'a' when $apr1$ or '1' when $1$
 } crypt_md5_salt;
 
 typedef struct {
-	unsigned char length;
-	unsigned char v[PLAINTEXT_LENGTH];
+	uint8_t length;
+	uint8_t v[PLAINTEXT_LENGTH];
 } crypt_md5_password;
 
 typedef struct {
-	uint32_t v[4];		//128 bits
-} crypt_md5_hash;
+	  char cracked;
+} crypt_md5_crack;
 
 typedef struct __attribute__((__aligned__(4))){
 	uint8_t buffer[64];
@@ -43,8 +44,6 @@ typedef struct __attribute__((__aligned_
 static const char md5_salt_prefix[] = "$1$";
 static const char apr1_salt_prefix[] = "$apr1$";
 
-#define address(j,idx) 			(((j)*KEYS_PER_CRYPT)+(idx))
-
 #define ROTATE_LEFT(x, s) ((x << s) | (x >> (32 - s)))
 
 #define F(x, y, z) ((z) ^ ((x) & ((y) ^ (z))))
diff -urpN magnum-jumbo/src/cuda_cryptmd5_fmt.c magnum-jumbo_cudatsfix//src/cuda_cryptmd5_fmt.c
--- magnum-jumbo/src/cuda_cryptmd5_fmt.c	2012-06-24 21:19:23.801544372 +0000
+++ magnum-jumbo_cudatsfix//src/cuda_cryptmd5_fmt.c	2012-06-24 21:38:20.287577288 +0000
@@ -1,14 +1,14 @@
 /*
-* This software is Copyright (c) 2011 Lukas Odzioba <lukas dot odzioba at gmail dot 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.
 */
 #include <string.h>
 #include <unistd.h>
 #include "arch.h"
-#include "formats.h"
-#include "common.h"
 #include "misc.h"
+#include "common.h"
+#include "formats.h"
 #include "cuda_common.h"
 #include "cuda_cryptmd5.h"
 
@@ -21,93 +21,98 @@
 #define BENCHMARK_LENGTH	-1
 
 #define BINARY_SIZE		16
-#define SALT_SIZE		(8+1)
+#define SALT_SIZE		(sizeof(crypt_md5_salt))
 #define MIN_KEYS_PER_CRYPT	KEYS_PER_CRYPT
 #define MAX_KEYS_PER_CRYPT	KEYS_PER_CRYPT
 
-void md5_crypt_gpu(crypt_md5_password *,uint32_t *, crypt_md5_salt *);
+void md5_crypt_gpu(crypt_md5_password *, crypt_md5_crack *, crypt_md5_salt *);
 
-static crypt_md5_password *inbuffer;//[MAX_KEYS_PER_CRYPT];			/** plaintext ciphertexts **/
-static uint32_t *outbuffer;//[MAX_KEYS_PER_CRYPT*4];			/** calculated hashes **/
-static crypt_md5_salt host_salt;					/** salt **/
+static crypt_md5_password *inbuffer;			/** plaintext ciphertexts **/
+static crypt_md5_crack *outbuffer;			/** cracked or no **/
+static crypt_md5_salt host_salt;			/** salt **/
+static int any_cracked;
 
+//#define CUDA_DEBUG
 
-//tests are unified for 8+8 length
 static struct fmt_tests tests[] = {
-/*	   {"$1$Btiy90iG$bGn4vzF3g1rIVGZ5odGIp/","qwerty"},
-	   {"$1$salt$c813W/s478KCzR0NnHx7j0","qwerty"},
-	   {"$1$salt$8LO.EVfsTf.HATV1Bd0ZP/","john"},
-	   {"$1$salt$TelRRxWBCxlpXmgAeB82R/","openwall"},
-	   {"$1$salt$l9PzDiECW83MOIMFTRL4Y1","summerofcode"},
-	   {"$1$salt$wZ2yVsplRoPoD7IfTvRsa0","IamMD5"},
-	   {"$1$saltstri$9S4.PyBpUZBRZw6ZsmFQE/","john"},
-	   {"$1$saltstring$YmP55hH3qcHg2cCffyxrq/","ala"},
-*/
-//         {"$1$salt1234$mdji1uBBCWZ5m2mIWKvLW.","a"},
-//         {"$1$salt1234$/JUvhIWHD.csWSCPvr7po0","ab"},
-//         {"$1$salt1234$GrxHg1bgkN2HB5CRCdrmF.","abc"},
-//         {"$1$salt1234$iZuyvTkrucWx8kVn5BN4M/","abcd"},
-//         {"$1$salt1234$wn0RbuDtbJlD1Q.X7.9wG/","abcde"},
-
-//         {"$1$salt1234$lzB83HS4FjzbcD4yMcjl01","abcdef"},
-//          {"$1$salt1234$bklJHN73KS04Kh6j6qPnr.","abcdefg"}, 
-	{"$1$salt1234$u4RMKGXG2b/Ud2rFmhqi70", "abcdefgh"},	//saltlen=8,passlen=8
-//         {"$1$salt1234$QjP48HUerU7aUYc/aJnre1","abcdefghi"},
-//         {"$1$salt1234$9jmu9ldi9vNw.XDO3TahR.","abcdefghij"},
-
-//         {"$1$salt1234$d3.LnlDWfkTIej5Ef1sCU/","abcdefghijk"},
-//         {"$1$salt1234$pDV0xEgZR14EpQMmhZ6Hg0","abcdefghijkl"},
-//         {"$1$salt1234$WumpbolX2y45Dlv0.A1Mj1","abcdefghijklm"},
-//         {"$1$salt1234$FXBreA27b7N7diemBGn5I1","abcdefghijklmn"},
-//         {"$1$salt1234$8d5IPIbTd7J/WNEG4b4cl.","abcdefghijklmno"},
-
-	//tests from korelogic2010 contest
-/*	   {"$1$bn6UVs3/$S6CQRLhmenR8OmVp3Jm5p0","sparky"},
-	   {"$1$qRiPuG5Z$pLLczmBnwEOD75Vb7YZLg1","walter"},
-	   {"$1$E.qsK.Hy$.eX0H6arTHaGOIFkf6o.a.","heaven"},
-	   {"$1$Hul2mrWs$.NGCgz3fBGDyG7RMGJAdM0","bananas"},
-	   {"$1$1l88Y.UV$swt2d0SPMrBPkdAD8RwSj0","horses"},
-	   {"$1$DiHrL6V7$fCVDD1GEAKB.BjAgJL1ZX0","maddie"},
-	   {"$1$7fpfV7kr$7LgF64DGPtHPktVKdLM490","bitch1"},
-	   {"$1$VKjk2PJc$5wbrtc9oa8kdEO/ocyi06/","crystal"},
-	   {"$1$S66DxkFm$kG.QfeHNLifEDTDmf4pzJ/","claudia"},
-	   {"$1$T2JMeEYj$Y.wDzFvyb9nlH1EiSCI3M/","august"}, 
-	 
-	   //tests from MD5_fmt.c
-	   {"$1$12345678$aIccj83HRDBo6ux1bVx7D1", "0123456789ABCDE"},
-	   {"$apr1$Q6ZYh...$RV6ft2bZ8j.NGrxLYaJt9.", "test"},
-	   {"$1$12345678$f8QoJuo0DpBRfQSD0vglc1", "12345678"},
-	   {"$1$$qRPK7m23GJusamGpoGLby/", ""},
-	   {"$apr1$a2Jqm...$grFrwEgiQleDr0zR4Jx1b.", "15 chars is max"},
-	   {"$1$$AuJCr07mI7DSew03TmBIv/", "no salt"},
-	   {"$1$`!@#%^&*$E6hD76/pKTS8qToBCkux30", "invalid salt"},
-	   {"$1$12345678$xek.CpjQUVgdf/P2N9KQf/", ""},
-	   {"$1$1234$BdIMOAWFOV2AQlLsrN/Sw.", "1234"},
-	   {"$apr1$rBXqc...$NlXxN9myBOk95T0AyLAsJ0", "john"},
-	   {"$apr1$Grpld/..$qp5GyjwM2dnA5Cdej9b411", "the"},
-	   {"$apr1$GBx.D/..$yfVeeYFCIiEXInfRhBRpy/", "ripper"},
+	{"$1$Btiy90iG$bGn4vzF3g1rIVGZ5odGIp/", "qwerty"},
+	/*{"$1$salt$c813W/s478KCzR0NnHx7j0", "qwerty"},
+	{"$1$salt$8LO.EVfsTf.HATV1Bd0ZP/", "john"},
+	{"$1$salt$TelRRxWBCxlpXmgAeB82R/", "openwall"},
+	{"$1$salt$l9PzDiECW83MOIMFTRL4Y1", "summerofcode"},
+	{"$1$salt$wZ2yVsplRoPoD7IfTvRsa0", "IamMD5"},
+	{"$1$saltstri$9S4.PyBpUZBRZw6ZsmFQE/", "john"},
+	{"$1$saltstring$YmP55hH3qcHg2cCffyxrq/", "ala"},
+
+	{"$1$salt1234$mdji1uBBCWZ5m2mIWKvLW.", "a"},
+	{"$1$salt1234$/JUvhIWHD.csWSCPvr7po0", "ab"},
+	{"$1$salt1234$GrxHg1bgkN2HB5CRCdrmF.", "abc"},
+	{"$1$salt1234$iZuyvTkrucWx8kVn5BN4M/", "abcd"},
+	{"$1$salt1234$wn0RbuDtbJlD1Q.X7.9wG/", "abcde"},
+
+	{"$1$salt1234$lzB83HS4FjzbcD4yMcjl01", "abcdef"},
+	{"$1$salt1234$bklJHN73KS04Kh6j6qPnr.", "abcdefg"},
+	{"$1$salt1234$u4RMKGXG2b/Ud2rFmhqi70", "abcdefgh"},
+	{"$1$salt1234$QjP48HUerU7aUYc/aJnre1", "abcdefghi"},
+	{"$1$salt1234$9jmu9ldi9vNw.XDO3TahR.", "abcdefghij"},
+
+	{"$1$salt1234$d3.LnlDWfkTIej5Ef1sCU/", "abcdefghijk"},
+	{"$1$salt1234$pDV0xEgZR14EpQMmhZ6Hg0", "abcdefghijkl"},
+	{"$1$salt1234$WumpbolX2y45Dlv0.A1Mj1", "abcdefghijklm"},
+	{"$1$salt1234$FXBreA27b7N7diemBGn5I1", "abcdefghijklmn"},
+	{"$1$salt1234$8d5IPIbTd7J/WNEG4b4cl.", "abcdefghijklmno"},
+
+	///tests from korelogic2010 contest
+	{"$1$bn6UVs3/$S6CQRLhmenR8OmVp3Jm5p0", "sparky"},
+	{"$1$qRiPuG5Z$pLLczmBnwEOD75Vb7YZLg1", "walter"},
+	{"$1$E.qsK.Hy$.eX0H6arTHaGOIFkf6o.a.", "heaven"},
+	{"$1$Hul2mrWs$.NGCgz3fBGDyG7RMGJAdM0", "bananas"},
+	{"$1$1l88Y.UV$swt2d0SPMrBPkdAD8RwSj0", "horses"},
+	{"$1$DiHrL6V7$fCVDD1GEAKB.BjAgJL1ZX0", "maddie"},
+	{"$1$7fpfV7kr$7LgF64DGPtHPktVKdLM490", "bitch1"},
+	{"$1$VKjk2PJc$5wbrtc9oa8kdEO/ocyi06/", "crystal"},
+	{"$1$S66DxkFm$kG.QfeHNLifEDTDmf4pzJ/", "claudia"},
+	{"$1$T2JMeEYj$Y.wDzFvyb9nlH1EiSCI3M/", "august"},
+
+	///tests from MD5_fmt.c
+	{"$1$12345678$aIccj83HRDBo6ux1bVx7D1", "0123456789ABCDE"},
+	{"$apr1$Q6ZYh...$RV6ft2bZ8j.NGrxLYaJt9.", "test"},
+	{"$1$12345678$f8QoJuo0DpBRfQSD0vglc1", "12345678"},
+	{"$1$$qRPK7m23GJusamGpoGLby/", ""},
+	{"$apr1$a2Jqm...$grFrwEgiQleDr0zR4Jx1b.", "15 chars is max"},
+	{"$1$$AuJCr07mI7DSew03TmBIv/", "no salt"},
+	{"$1$`!@#%^&*$E6hD76/pKTS8qToBCkux30", "invalid salt"},
+	{"$1$12345678$xek.CpjQUVgdf/P2N9KQf/", ""},
+	{"$1$1234$BdIMOAWFOV2AQlLsrN/Sw.", "1234"},
+	{"$apr1$rBXqc...$NlXxN9myBOk95T0AyLAsJ0", "john"},
+	{"$apr1$Grpld/..$qp5GyjwM2dnA5Cdej9b411", "the"},
+	{"$apr1$GBx.D/..$yfVeeYFCIiEXInfRhBRpy/", "ripper"},
 */
 	{NULL}
 };
 
 static void cleanup()
 {
- free(inbuffer);
- free(outbuffer);
+	free(inbuffer);
+	free(outbuffer);
 }
 
 static void init(struct fmt_main *pFmt)
-{  
-  //Alocate memory for hashes and passwords
-  inbuffer=(crypt_md5_password*)malloc(sizeof(crypt_md5_password)*MAX_KEYS_PER_CRYPT);
-  outbuffer=(uint32_t*)malloc(sizeof(uint32_t)*MAX_KEYS_PER_CRYPT*4);
-  check_mem_allocation(inbuffer,outbuffer);
-  atexit(cleanup);
-  //Initialize CUDA
-  cuda_init(gpu_id);
+{
+	///Alocate memory for hashes and passwords
+	inbuffer =
+	    (crypt_md5_password *) calloc(MAX_KEYS_PER_CRYPT,
+	    sizeof(crypt_md5_password));
+	outbuffer =
+	    (crypt_md5_crack *) calloc(MAX_KEYS_PER_CRYPT,
+	    sizeof(crypt_md5_crack));
+	check_mem_allocation(inbuffer, outbuffer);
+	atexit(cleanup);
+	///Initialize CUDA
+	cuda_init(gpu_id);
 }
 
-static int valid(char *ciphertext,struct fmt_main *pFmt)
+static int valid(char *ciphertext, struct fmt_main *pFmt)
 {
 	uint8_t i, len = strlen(ciphertext), prefix = 0;
 
@@ -118,7 +123,6 @@ static int valid(char *ciphertext,struct
 		prefix |= 2;
 	if (prefix == 0)
 		return 0;
-
 	char *p = strrchr(ciphertext, '$');
 	if (p == NULL)
 		return 0;
@@ -155,7 +159,6 @@ static void to_binary(char *crypt, char
       alt[B1]=b1;\
       alt[B0]=b2;\
   }
-
 	_24bit_from_b64(0, 0, 6, 12);
 	_24bit_from_b64(4, 1, 7, 13);
 	_24bit_from_b64(8, 2, 8, 14);
@@ -177,72 +180,50 @@ static void *binary(char *ciphertext)
 
 static void *salt(char *ciphertext)
 {
-	static uint8_t ret[SALT_SIZE];
-	memset(ret, 0, SALT_SIZE);
-	uint8_t i, *pos = (uint8_t *) ciphertext, *dest = ret, *end;
-
+#ifdef CUDA_DEBUG
+	printf("salt(%s)\n", ciphertext);
+#endif
+	static crypt_md5_salt ret;
+	uint8_t i, *pos = (uint8_t *) ciphertext, *end;
+	char *dest = ret.salt;
 	if (strncmp(ciphertext, md5_salt_prefix, strlen(md5_salt_prefix)) == 0) {
 		pos += strlen(md5_salt_prefix);
-		ret[8] = '1';
+		ret.prefix = '1';
 	}
 	if (strncmp(ciphertext, apr1_salt_prefix,
 		strlen(apr1_salt_prefix)) == 0) {
 		pos += strlen(apr1_salt_prefix);
-		ret[8] = 'a';
+		ret.prefix = 'a';
 	}
 	end = pos;
 	for (i = 0; i < 8 && *end != '$'; i++, end++);
 	while (pos != end)
 		*dest++ = *pos++;
-	return (void *) ret;
-}
-
-static int binary_hash_0(void *binary)
-{
-	return (((ARCH_WORD_32 *) binary)[0] & 0xf);
-}
-
-static int binary_hash_1(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0xff;
-}
-
-static int binary_hash_2(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0xfff;
-}
-
-static int binary_hash_3(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0xffff;
-}
-
-static int binary_hash_4(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0xfffff;
-}
-
-static int binary_hash_5(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0xffffff;
+	ret.length = i;	
+	char *p = strrchr(ciphertext, '$') + 1;
+	to_binary(p,(char*) ret.hash);
+#ifdef CUDA_DEBUG
+	puts("salted:");
+	uint32_t *t=ret.hash;
+	for(i=0;i<4;i++)
+	  printf("%08x ",t[i]);
+	puts("");
+#endif
+	return (void *) &ret;
 }
 
-static int binary_hash_6(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0x7ffffff;
-}
 static void set_salt(void *salt)
 {
-	uint8_t *s = salt;
-	uint8_t len;
-	for (len = 0; len < 8 && s[len]; len++);
-	host_salt.saltlen = len;
-	memcpy(host_salt.salt, s, host_salt.saltlen);
-	host_salt.prefix = s[8];
+	memcpy(&host_salt, salt, sizeof(crypt_md5_salt));
+	any_cracked = 0;
 }
 
 static void set_key(char *key, int index)
 {
+  
+#ifdef CUDA_DEBUG
+	printf("set_key(%d,%s)\n", index, key);
+#endif
 	uint32_t len = strlen(key);
 	inbuffer[index].length = len;
 	memcpy((char *) inbuffer[index].v, key, len);
@@ -258,65 +239,34 @@ static char *get_key(int index)
 
 static void crypt_all(int count)
 {
+	int i;
+	if (any_cracked) {
+		memset(outbuffer, 0, sizeof(crypt_md5_crack) * KEYS_PER_CRYPT);
+		any_cracked = 0;
+	}
 	md5_crypt_gpu(inbuffer, outbuffer, &host_salt);
-}
-
-static int get_hash_0(int index)
-{
-	return outbuffer[address(0, index)] & 0xf;
-}
-
-static int get_hash_1(int index)
-{
-	return outbuffer[address(0, index)] & 0xff;
-}
-
-static int get_hash_2(int index)
-{
-	return outbuffer[address(0, index)] & 0xfff;
-}
-
-static int get_hash_3(int index)
-{
-	return outbuffer[address(0, index)] & 0xffff;
-}
-
-static int get_hash_4(int index)
-{
-	return outbuffer[address(0, index)] & 0xfffff;
-}
-
-static int get_hash_5(int index)
-{
-	return outbuffer[address(0, index)] & 0xffffff;
-}
-
-static int get_hash_6(int index)
-{
-	return outbuffer[address(0, index)] & 0x7ffffff;
+	for (i = 0; i < count; i++) {
+		any_cracked|=outbuffer[i].cracked;
+	}
+#ifdef CUDA_DEBUG
+	printf("crypt_all(%d)\n", count);
+	printf("any_cracked=%d\n",any_cracked);
+#endif
 }
 
 static int cmp_all(void *binary, int count)
 {
-	uint32_t i, b = ((uint32_t *) binary)[0];
-	for (i = 0; i < count; i++)
-		if(b==outbuffer[address(0,i)])
-		return 1;
-	return 0;
+	return any_cracked;
 }
 
 static int cmp_one(void *binary, int index)
 {
-	uint32_t i, *t = (uint32_t *) binary;
-	for (i = 0; i < 4; i++)
-		if(t[i]!=outbuffer[address(i, index)])
-		  return 0;
-	return 1;
+	return outbuffer[index].cracked;
 }
 
-static int cmp_exact(char *source, int count)
+static int cmp_exact(char *source, int index)
 {
-	return 1;
+	return outbuffer[index].cracked;
 }
 
 struct fmt_main fmt_cuda_cryptmd5 = {
@@ -332,7 +282,8 @@ struct fmt_main fmt_cuda_cryptmd5 = {
 		    MIN_KEYS_PER_CRYPT,
 		    MAX_KEYS_PER_CRYPT,
 		    FMT_CASE | FMT_8_BIT,
-	    tests},
+		    tests
+	},
 	{
 		    init,
 		    fmt_default_prepare,
@@ -341,13 +292,7 @@ struct fmt_main fmt_cuda_cryptmd5 = {
 		    binary,
 		    salt,
 		    {
-				binary_hash_0,
-				binary_hash_1,
-				binary_hash_2,
-				binary_hash_3,
-				binary_hash_4,
-				binary_hash_5,
-				binary_hash_6
+			fmt_default_binary_hash
 		    },
 		    fmt_default_salt_hash,
 		    set_salt,
@@ -356,15 +301,10 @@ struct fmt_main fmt_cuda_cryptmd5 = {
 		    fmt_default_clear_keys,
 		    crypt_all,
 		    {
-				get_hash_0,
-				get_hash_1,
-				get_hash_2,
-				get_hash_3,
-			get_hash_4,
-			get_hash_5,
-			get_hash_6
+			fmt_default_get_hash
 		    },
 		    cmp_all,
 		    cmp_one,
-	    cmp_exact}
+		    cmp_exact
+	}
 };
diff -urpN magnum-jumbo/src/cuda_cryptsha256_fmt.c magnum-jumbo_cudatsfix//src/cuda_cryptsha256_fmt.c
--- magnum-jumbo/src/cuda_cryptsha256_fmt.c	2012-06-24 21:19:23.801544372 +0000
+++ magnum-jumbo_cudatsfix//src/cuda_cryptsha256_fmt.c	2012-06-24 21:37:12.717577327 +0000
@@ -90,7 +90,7 @@ static void cleanup()
 static void init(struct fmt_main *pFmt)
 {
   //Alocate memory for hashes and passwords
-  inbuffer=(crypt_sha256_password*)malloc(sizeof(crypt_sha256_password)*MAX_KEYS_PER_CRYPT);
+  inbuffer=(crypt_sha256_password*)calloc(MAX_KEYS_PER_CRYPT,sizeof(crypt_sha256_password));
   outbuffer=(uint32_t*)malloc(sizeof(uint32_t)*MAX_KEYS_PER_CRYPT*8);
   check_mem_allocation(inbuffer,outbuffer);
   atexit(cleanup);
diff -urpN magnum-jumbo/src/cuda_cryptsha512_fmt.c magnum-jumbo_cudatsfix//src/cuda_cryptsha512_fmt.c
--- magnum-jumbo/src/cuda_cryptsha512_fmt.c	2012-06-24 21:19:23.801544372 +0000
+++ magnum-jumbo_cudatsfix//src/cuda_cryptsha512_fmt.c	2012-06-24 21:40:38.010576847 +0000
@@ -70,7 +70,7 @@ static void cleanup()
 static void init(struct fmt_main *pFmt)
 {
   //Alocate memory for hashes and passwords
-  inbuffer=(crypt_sha512_password*)malloc(sizeof(crypt_sha512_password)*MAX_KEYS_PER_CRYPT);
+  inbuffer=(crypt_sha512_password*)calloc(MAX_KEYS_PER_CRYPT,sizeof(crypt_sha512_password));
   outbuffer=(crypt_sha512_hash*)malloc(sizeof(crypt_sha512_hash)*MAX_KEYS_PER_CRYPT);
   check_mem_allocation(inbuffer,outbuffer);
   atexit(cleanup);
diff -urpN magnum-jumbo/src/cuda_mscash2_fmt.c magnum-jumbo_cudatsfix//src/cuda_mscash2_fmt.c
--- magnum-jumbo/src/cuda_mscash2_fmt.c	2012-06-24 21:19:23.801544372 +0000
+++ magnum-jumbo_cudatsfix//src/cuda_mscash2_fmt.c	2012-06-24 21:42:21.280624486 +0000
@@ -46,7 +46,7 @@ static void cleanup()
 static void init(struct fmt_main *pFmt)
 {
   //Alocate memory for hashes and passwords
-  inbuffer=(mscash2_password*)malloc(sizeof(mscash2_password)*MAX_KEYS_PER_CRYPT);
+  inbuffer=(mscash2_password*)calloc(MAX_KEYS_PER_CRYPT,sizeof(mscash2_password));
   outbuffer=(mscash2_hash*)malloc(sizeof(mscash2_hash)*MAX_KEYS_PER_CRYPT);
   check_mem_allocation(inbuffer,outbuffer);
   atexit(cleanup);
diff -urpN magnum-jumbo/src/cuda_mscash_fmt.c magnum-jumbo_cudatsfix//src/cuda_mscash_fmt.c
--- magnum-jumbo/src/cuda_mscash_fmt.c	2012-06-24 21:19:23.801544372 +0000
+++ magnum-jumbo_cudatsfix//src/cuda_mscash_fmt.c	2012-06-24 21:46:22.036576915 +0000
@@ -40,7 +40,7 @@ static void cleanup()
 static void init(struct fmt_main *pFmt)
 {
   //Alocate memory for hashes and passwords
-  inbuffer=(mscash_password*)malloc(sizeof(mscash_password)*MAX_KEYS_PER_CRYPT);
+  inbuffer=(mscash_password*)calloc(MAX_KEYS_PER_CRYPT, sizeof(mscash_password)); 
   outbuffer=(mscash_hash*)malloc(sizeof(mscash_hash)*MAX_KEYS_PER_CRYPT);
   check_mem_allocation(inbuffer,outbuffer);
   atexit(cleanup);
diff -urpN magnum-jumbo/src/cuda_phpass.h magnum-jumbo_cudatsfix//src/cuda_phpass.h
--- magnum-jumbo/src/cuda_phpass.h	2012-06-24 21:19:23.801544372 +0000
+++ magnum-jumbo_cudatsfix//src/cuda_phpass.h	2012-06-20 03:06:51.000000000 +0000
@@ -1,6 +1,5 @@
 /*
-* This software is Copyright (c) 2011 Lukas Odzioba
-* <lukas dot odzioba at gmail dot 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.
 */
@@ -11,9 +10,10 @@
 #define uint32_t 			unsigned int
 
 #define ROTATE_LEFT(x, s) 		((x << s) | (x >> (32 - s)))
-#define BLOCKS 				126*3	//it must be always something*3
+#define BLOCKS 				126
 #define THREADS 			256
 #define KEYS_PER_CRYPT 			BLOCKS*THREADS
+#define SALT_SIZE			sizeof(phpass_salt)
 
 #define F(x, y, z)			((z) ^ ((x) & ((y) ^ (z))))
 #define G(x, y, z)			((y) ^ ((z) & ((x) ^ (y))))
@@ -70,12 +70,18 @@
 static char phpass_prefix[] = "$P$";
 
 typedef struct {
-	unsigned char v[15];
-	unsigned char length;
+	uint8_t v[15];
+	uint8_t length;
 } phpass_password;
 
 typedef struct {
-	uint32_t v[4];		///128bits for hash
-} phpass_hash;
+	uint8_t salt[8];
+	uint32_t hash[4];
+	uint32_t rounds;
+} phpass_salt;
+
+typedef struct {
+	uint8_t cracked;
+} phpass_crack;
 
 #endif
diff -urpN magnum-jumbo/src/cuda_phpass_fmt.c magnum-jumbo_cudatsfix//src/cuda_phpass_fmt.c
--- magnum-jumbo/src/cuda_phpass_fmt.c	2012-06-24 21:19:23.801544372 +0000
+++ magnum-jumbo_cudatsfix//src/cuda_phpass_fmt.c	2012-06-24 21:50:23.803452860 +0000
@@ -1,5 +1,5 @@
 /*
-* This software is Copyright (c) 2011 Lukas Odzioba <lukas dot odzioba at gmail dot 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.
 */
@@ -17,6 +17,7 @@
 #define ALGORITHM_NAME		"CUDA"
 
 #define BENCHMARK_COMMENT	" ($P$9 lengths 1 to 15)"
+
 #define BENCHMARK_LENGTH	-1
 
 #define PLAINTEXT_LENGTH	15
@@ -24,25 +25,18 @@
 #define BINARY_SIZE		16
 #define MD5_DIGEST_LENGTH 	16
 
-
-#define SALT_SIZE		8
-
 #define MIN_KEYS_PER_CRYPT	KEYS_PER_CRYPT
 #define MAX_KEYS_PER_CRYPT	KEYS_PER_CRYPT
 
-static unsigned char *inbuffer;//[MAX_KEYS_PER_CRYPT * sizeof(phpass_password)];			/** plaintext ciphertexts **/
-static uint32_t *outbuffer;//[MAX_KEYS_PER_CRYPT * 4];						/** calculated hashes **/
-
-static char currentsalt[SALT_SIZE];
-static char loopChar = '*';
-
-extern void mem_init(unsigned char *, uint32_t *, char *, char *, int);
-extern void mem_clear(void);
-extern void gpu_phpass(void);
+static unsigned char *inbuffer;				/** plaintext ciphertexts **/
+static phpass_crack *outbuffer;				/** calculated hashes **/
+static phpass_salt currentsalt;
+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$900000000jPBDh/JWJIyrF0.DmP7kT.", "ala"},
 	{"$P$900000000a94rg7R/nUK0icmALICKj1", "john"},
 	{"$P$900000001ahWiA6cMRZxkgUxj4x/In0", "john"},
 	{"$P$900000000m6YEJzWtTmNBBL4jypbHv1", "openwall"},
@@ -52,44 +46,46 @@ static struct fmt_tests tests[] = {
 	{"$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$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"},
 	{NULL}
 };
 
 static void cleanup()
 {
-  free(inbuffer);
-  free(outbuffer);
+	free(inbuffer);
+	free(outbuffer);
 }
 
 static void init(struct fmt_main *pFmt)
 {
-    //Alocate memory for hashes and passwords
-    inbuffer=(unsigned char*)malloc(MAX_KEYS_PER_CRYPT * sizeof(phpass_password)*sizeof(char));
-    outbuffer=(uint32_t *)malloc(MAX_KEYS_PER_CRYPT*4*sizeof(uint32_t));
-    check_mem_allocation(inbuffer,outbuffer);
-    atexit(cleanup);
-    //Initialize CUDA
-    cuda_init(gpu_id);
+	///Alocate memory for hashes and passwords
+	inbuffer =
+	    (uint8_t *) calloc(MAX_KEYS_PER_CRYPT, sizeof(phpass_password));
+	outbuffer =
+	    (phpass_crack *) calloc(MAX_KEYS_PER_CRYPT, sizeof(phpass_crack));
+	check_mem_allocation(inbuffer, outbuffer);
+	atexit(cleanup);
+	///Initialize CUDA
+	cuda_init(gpu_id);
 }
 
-static int valid(char *ciphertext,struct fmt_main *pFmt)
+static int valid(char *ciphertext, struct fmt_main *pFmt)
 {
 	uint32_t i, count_log2;
 
@@ -109,85 +105,53 @@ static int valid(char *ciphertext,struct
 	return 1;
 };
 
-//code from historical JtR phpass patch
-static void *binary(char *ciphertext)
+///code from historical JtR phpass patch
+static void pbinary(char *ciphertext, unsigned char *out)
 {
-	static unsigned char b[BINARY_SIZE];
-	memset(b, 0, BINARY_SIZE);
+	memset(out, 0, BINARY_SIZE);
 	int i, bidx = 0;
 	unsigned sixbits;
 	char *pos = &ciphertext[3 + 1 + 8];
 
 	for (i = 0; i < 5; i++) {
 		sixbits = atoi64[ARCH_INDEX(*pos++)];
-		b[bidx] = sixbits;
+		out[bidx] = sixbits;
 		sixbits = atoi64[ARCH_INDEX(*pos++)];
-		b[bidx++] |= (sixbits << 6);
+		out[bidx++] |= (sixbits << 6);
 		sixbits >>= 2;
-		b[bidx] = sixbits;
+		out[bidx] = sixbits;
 		sixbits = atoi64[ARCH_INDEX(*pos++)];
-		b[bidx++] |= (sixbits << 4);
+		out[bidx++] |= (sixbits << 4);
 		sixbits >>= 4;
-		b[bidx] = sixbits;
+		out[bidx] = sixbits;
 		sixbits = atoi64[ARCH_INDEX(*pos++)];
-		b[bidx++] |= (sixbits << 2);
+		out[bidx++] |= (sixbits << 2);
 	}
 	sixbits = atoi64[ARCH_INDEX(*pos++)];
-	b[bidx] = sixbits;
+	out[bidx] = sixbits;
 	sixbits = atoi64[ARCH_INDEX(*pos++)];
-	b[bidx] |= (sixbits << 6);
-	return (void *) b;
-}
-
-static void *salt(char *ciphertext)
-{
-	static unsigned char salt[SALT_SIZE + 2];
-	memcpy(salt, &ciphertext[4], 8);
-	salt[8] = ciphertext[3];
-	salt[9] = 0;
-	return salt;
-}
-
-static int binary_hash_0(void *binary)
-{
-	return (((ARCH_WORD_32 *) binary)[0] & 0xf);
-}
-
-static int binary_hash_1(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0xff;
-}
-
-static int binary_hash_2(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0xfff;
-}
-
-static int binary_hash_3(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0xffff;
-}
-
-static int binary_hash_4(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0xfffff;
+	out[bidx] |= (sixbits << 6);
 }
 
-static int binary_hash_5(void *binary)
+static void *binary(char *ciphertext)
 {
-	return ((ARCH_WORD_32 *) binary)[0] & 0xffffff;
+	static unsigned char b[BINARY_SIZE];
+	pbinary(ciphertext, b);
+	return (void *) b;
 }
 
-static int binary_hash_6(void *binary)
+static void *salt(char *ciphertext)
 {
-	return ((ARCH_WORD_32 *) binary)[0] & 0x7ffffff;
+	static phpass_salt salt;
+	salt.rounds = 1 << atoi64[ARCH_INDEX(ciphertext[3])];
+	memcpy(salt.salt, &ciphertext[4], 8);
+	pbinary(ciphertext, (unsigned char*)salt.hash);
+	return &salt;
 }
 
 static void set_salt(void *salt)
 {
-	unsigned char *csalt = salt;
-	memcpy(currentsalt,csalt,8);
-	loopChar = csalt[8];
+	memcpy(&currentsalt, salt, SALT_SIZE);
 }
 
 static void set_key(char *key, int index)
@@ -204,82 +168,36 @@ static char *get_key(int index)
 	int i;
 	for (i = 0; i < PLAINTEXT_LENGTH; i++)
 		r[i] = inbuffer[address(i, index)];
-	r[inbuffer[address(15,index)]] = '\0';
+	r[inbuffer[address(15, index)]] = '\0';
 	return r;
 }
 
 static void crypt_all(int count)
 {
-	char setting[40];
-	strcpy(setting, phpass_prefix);
-	setting[3] = loopChar;
-	int count_log2 = 0;
-	count_log2 = atoi64[ARCH_INDEX(setting[3])];
-	strcpy(setting + 4, currentsalt);
-	mem_init(inbuffer, outbuffer, setting, itoa64, count_log2);
-	gpu_phpass();
-	mem_clear();
-}
-
-static int get_hash_0(int index)
-{
-	return outbuffer[address(0, index)] & 0xf;
-}
-
-static int get_hash_1(int index)
-{
-	return outbuffer[address(0, index)] & 0xff;
-}
-
-static int get_hash_2(int index)
-{
-	return outbuffer[address(0, index)] & 0xfff;
-}
-
-static int get_hash_3(int index)
-{
-	return outbuffer[address(0, index)] & 0xffff;
-}
-
-static int get_hash_4(int index)
-{
-	return outbuffer[address(0, index)] & 0xfffff;
-}
-
-static int get_hash_5(int index)
-{
-	return outbuffer[address(0, index)] & 0xffffff;
-}
-
-static int get_hash_6(int index)
-{
-	return outbuffer[address(0, index)] & 0x7ffffff;
+	int i;
+	if (any_cracked) {
+		memset(outbuffer, 0, sizeof(phpass_crack) * KEYS_PER_CRYPT);
+		any_cracked = 0;
+	}
+	gpu_phpass(inbuffer, &currentsalt, outbuffer);
+	for (i = 0; i < count; i++) {
+		any_cracked |= outbuffer[i].cracked;
+	}
 }
 
-
 static int cmp_all(void *binary, int count)
 {
-	uint32_t i;
-	uint32_t b = ((uint32_t *) binary)[0];
-	for (i = 0; i < count; i++)
-		if (b == outbuffer[address(0, i)])
-			return 1;
-	return 0;
+	return any_cracked;
 }
 
 static int cmp_one(void *binary, int index)
 {
-	int i;
-	uint32_t *t = (uint32_t *) binary;
-	for (i = 0; i < 4; i++)
-		if (t[i] != outbuffer[address(i, index)])
-			return 0;
-	return 1;
+	return outbuffer[index].cracked;
 }
 
-static int cmp_exact(char *source, int count)
+static int cmp_exact(char *source, int index)
 {
-	return 1;
+	return outbuffer[index].cracked;
 }
 
 struct fmt_main fmt_cuda_phpass = {
@@ -291,11 +209,12 @@ struct fmt_main fmt_cuda_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,
-	    tests},
+		    tests
+	},
 	{
 		    init,
 		    fmt_default_prepare,
@@ -304,13 +223,8 @@ struct fmt_main fmt_cuda_phpass = {
 		    binary,
 		    salt,
 		    {
-				binary_hash_0,
-				binary_hash_1,
-				binary_hash_2,
-				binary_hash_3,
-			binary_hash_4,
-		    binary_hash_5,
-		    binary_hash_6},
+			fmt_default_binary_hash
+		    },
 		    fmt_default_salt_hash,
 		    set_salt,
 		    set_key,
@@ -318,14 +232,10 @@ struct fmt_main fmt_cuda_phpass = {
 		    fmt_default_clear_keys,
 		    crypt_all,
 		    {
-				get_hash_0,
-				get_hash_1,
-				get_hash_2,
-				get_hash_3,
-			get_hash_4,
-		    get_hash_5,
-		    get_hash_6},
+			fmt_default_get_hash
+		    },
 		    cmp_all,
 		    cmp_one,
-	    cmp_exact}
+		    cmp_exact
+	}
 };