Add binary_hash() functions to phpass-cuda and md5crypt-cuda formats.

src/cuda/cryptmd5.cu

@@ -131,7 +131,7 @@ __device__ void md5_digest(md5_ctx * ctx, uint32_t * result,
 }
 
 
-__device__ void md5crypt(const char *gpass, size_t keysize, char *result)
+__device__ void md5crypt(const char *gpass, size_t keysize, unsigned int *result)
 {
 
 	uint32_t i;
@@ -202,12 +202,10 @@ __device__ void md5crypt(const char *gpass, size_t keysize, char *result)
 			ctx_update(&ctx, pass, pass_len, &ctx_buflen);
 		md5_digest(&ctx, alt_result[threadIdx.x], &ctx_buflen);
 	}
-	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;
+	result[0] = alt_result[threadIdx.x][0];
+	result[1] = alt_result[threadIdx.x][1];
+	result[2] = alt_result[threadIdx.x][2];
+	result[3] = alt_result[threadIdx.x][3];
 }
 
 
@@ -216,7 +214,7 @@ __global__ void kernel_crypt_r(crypt_md5_password * inbuffer,
 {
 	uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
 	md5crypt((char *) inbuffer[idx].v, inbuffer[idx].length,
-	    &outbuffer[idx].cracked);
+	    outbuffer[idx].hash);
 }
 
 __host__ void md5_crypt_gpu(crypt_md5_password * inbuffer, uint32_t * outbuffer,

src/cuda/phpass.cu

@@ -266,11 +266,8 @@ __global__ void kernel_phpass(unsigned char *password, phpass_crack * data_out)
 		x3 = d + 0x10325476;
 
 	} while (--count);
-
-	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;
+	data_out[idx].hash[0] = x0;
+	data_out[idx].hash[1] = x1;
+	data_out[idx].hash[2] = x2;
+	data_out[idx].hash[3] = x3;
 }

src/cuda_cryptmd5.h

@@ -22,7 +22,6 @@
 #define PLAINTEXT_LENGTH	15
 
 typedef struct {
-	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$
@@ -34,7 +33,7 @@ typedef struct {
 } crypt_md5_password;
 
 typedef struct {
-	  char cracked;
+	uint32_t hash[4];
 } crypt_md5_crack;
 
 typedef struct __attribute__((__aligned__(4))){

src/cuda_cryptmd5_fmt.c

@@ -31,7 +31,6 @@ void md5_crypt_gpu(crypt_md5_password *, crypt_md5_crack *,
 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
 
@@ -191,7 +190,7 @@ static void *salt(char *ciphertext)
 #endif
 	static crypt_md5_salt ret;
 	uint8_t i, *pos = (uint8_t *) ciphertext, *end;
-	char *p,*dest = ret.salt;
+	char *dest = ret.salt;
 	if (strncmp(ciphertext, md5_salt_prefix, strlen(md5_salt_prefix)) == 0) {
 		pos += strlen(md5_salt_prefix);
 		ret.prefix = '1';
@@ -206,15 +205,6 @@ static void *salt(char *ciphertext)
 	while (pos != end)
 		*dest++ = *pos++;
 	ret.length = i;
-	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;
 }
 
@@ -244,39 +234,43 @@ static char *get_key(int index)
 
 static int crypt_all(int *pcount, struct db_salt *salt)
 {
-	int count = *pcount;
-	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, count);
-	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
-	return count;
+	md5_crypt_gpu(inbuffer, outbuffer, &host_salt, *pcount);
+	return *pcount;
 }
 
 static int cmp_all(void *binary, int count)
 {
-	return any_cracked;
+	int i;
+	unsigned int *b32 = (unsigned int *)binary;
+	for(i=0; i < count; i++)
+		if(outbuffer[i].hash[0] == b32[0])
+			return 1;
+	return 0;
 }
 
 static int cmp_one(void *binary, int index)
 {
-	return outbuffer[index].cracked;
+	int i;
+	unsigned int *b32 = (unsigned int *)binary;
+	for(i=0; i < 4; i++)
+		if(outbuffer[index].hash[i] != b32[i])
+			return 0;
+	return 1;
 }
 
 static int cmp_exact(char *source, int index)
 {
-	return outbuffer[index].cracked;
+	return 1;
 }
 
+static int get_hash_0(int index) { return outbuffer[index].hash[0] & 0xf; }
+static int get_hash_1(int index) { return outbuffer[index].hash[0] & 0xff; }
+static int get_hash_2(int index) { return outbuffer[index].hash[0] & 0xfff; }
+static int get_hash_3(int index) { return outbuffer[index].hash[0] & 0xffff; }
+static int get_hash_4(int index) { return outbuffer[index].hash[0] & 0xfffff; }
+static int get_hash_5(int index) { return outbuffer[index].hash[0] & 0xffffff; }
+static int get_hash_6(int index) { return outbuffer[index].hash[0] & 0x7ffffff; }
+
 struct fmt_main fmt_cuda_cryptmd5 = {
 	{
 		FORMAT_LABEL,
@@ -304,7 +298,13 @@ struct fmt_main fmt_cuda_cryptmd5 = {
 		salt,
 		fmt_default_source,
 		{
-			fmt_default_binary_hash
+			fmt_default_binary_hash_0,
+			fmt_default_binary_hash_1,
+			fmt_default_binary_hash_2,
+			fmt_default_binary_hash_3,
+			fmt_default_binary_hash_4,
+			fmt_default_binary_hash_5,
+			fmt_default_binary_hash_6
 		},
 		fmt_default_salt_hash,
 		set_salt,
@@ -313,7 +313,13 @@ struct fmt_main fmt_cuda_cryptmd5 = {
 		fmt_default_clear_keys,
 		crypt_all,
 		{
-			fmt_default_get_hash
+			get_hash_0,
+			get_hash_1,
+			get_hash_2,
+			get_hash_3,
+			get_hash_4,
+			get_hash_5,
+			get_hash_6
 		},
 		cmp_all,
 		cmp_one,

src/cuda_phpass.h

@@ -73,12 +73,11 @@ typedef struct {
 
 typedef struct {
 	uint8_t salt[8];
-	uint32_t hash[4];
 	uint32_t rounds;
 } phpass_salt;
 
 typedef struct {
-	uint8_t cracked;
+	uint32_t hash[4];
 } phpass_crack;
 
 #endif

src/cuda_phpass_fmt.c

@@ -31,7 +31,6 @@
 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 *, int count);
 
@@ -165,7 +164,6 @@ static void *salt(char *ciphertext)
 	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;
 }
 
@@ -194,35 +192,44 @@ static char *get_key(int index)
 
 static int crypt_all(int *pcount, struct db_salt *salt)
 {
-	int count = *pcount;
-	int i;
-
-	if (any_cracked) {
-		memset(outbuffer, 0, sizeof(phpass_crack) * KEYS_PER_CRYPT);
-		any_cracked = 0;
-	}
-	gpu_phpass(inbuffer, &currentsalt, outbuffer, count);
-	for (i = 0; i < count; i++) {
-		any_cracked |= outbuffer[i].cracked;
-	}
-	return count;
+	memset(outbuffer, 0, sizeof(phpass_crack) * KEYS_PER_CRYPT);
+	gpu_phpass(inbuffer, &currentsalt, outbuffer, *pcount);
+	return *pcount;
 }
 
 static int cmp_all(void *binary, int count)
 {
-	return any_cracked;
+	int i;
+	unsigned int *b32 = (unsigned int *)binary;
+	for(i=0; i < count; i++)
+		if(outbuffer[i].hash[0] == b32[0])
+			return 1;
+	return 0;
 }
 
 static int cmp_one(void *binary, int index)
 {
-	return outbuffer[index].cracked;
+	int i;
+	unsigned int *b32 = (unsigned int *)binary;
+	for(i=0; i < 4; i++)
+		if(outbuffer[index].hash[i] != b32[i])
+			return 0;
+	return 1;
 }
 
 static int cmp_exact(char *source, int index)
 {
-	return outbuffer[index].cracked;
+	return 1;
 }
 
+static int get_hash_0(int index) { return outbuffer[index].hash[0] & 0xf; }
+static int get_hash_1(int index) { return outbuffer[index].hash[0] & 0xff; }
+static int get_hash_2(int index) { return outbuffer[index].hash[0] & 0xfff; }
+static int get_hash_3(int index) { return outbuffer[index].hash[0] & 0xffff; }
+static int get_hash_4(int index) { return outbuffer[index].hash[0] & 0xfffff; }
+static int get_hash_5(int index) { return outbuffer[index].hash[0] & 0xffffff; }
+static int get_hash_6(int index) { return outbuffer[index].hash[0] & 0x7ffffff; }
+
 struct fmt_main fmt_cuda_phpass = {
 	{
 		FORMAT_LABEL,
@@ -250,7 +257,13 @@ struct fmt_main fmt_cuda_phpass = {
 		salt,
 		fmt_default_source,
 		{
-			fmt_default_binary_hash
+			fmt_default_binary_hash_0,
+			fmt_default_binary_hash_1,
+			fmt_default_binary_hash_2,
+			fmt_default_binary_hash_3,
+			fmt_default_binary_hash_4,
+			fmt_default_binary_hash_5,
+			fmt_default_binary_hash_6
 		},
 		fmt_default_salt_hash,
 		set_salt,
@@ -259,7 +272,13 @@ struct fmt_main fmt_cuda_phpass = {
 		fmt_default_clear_keys,
 		crypt_all,
 		{
-			fmt_default_get_hash
+			get_hash_0,
+			get_hash_1,
+			get_hash_2,
+			get_hash_3,
+			get_hash_4,
+			get_hash_5,
+			get_hash_6
 		},
 		cmp_all,
 		cmp_one,