diff -urpN jumbo.orig/src/cuda/cryptmd5.cu jumbo/src/cuda/cryptmd5.cu --- jumbo.orig/src/cuda/cryptmd5.cu 2013-08-05 23:59:46.179732854 +0000 +++ jumbo/src/cuda/cryptmd5.cu 2013-08-06 21:53:17.556908034 +0000 @@ -131,7 +131,7 @@ __device__ void md5_digest(md5_ctx * ctx } -__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 *gpa 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 { 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, diff -urpN jumbo.orig/src/cuda/phpass.cu jumbo/src/cuda/phpass.cu --- jumbo.orig/src/cuda/phpass.cu 2013-08-05 23:59:46.179732854 +0000 +++ jumbo/src/cuda/phpass.cu 2013-08-06 00:29:16.597042941 +0000 @@ -266,11 +266,8 @@ __global__ void kernel_phpass(unsigned c 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; } diff -urpN jumbo.orig/src/cuda_cryptmd5_fmt.c jumbo/src/cuda_cryptmd5_fmt.c --- jumbo.orig/src/cuda_cryptmd5_fmt.c 2013-08-05 23:59:46.159732736 +0000 +++ jumbo/src/cuda_cryptmd5_fmt.c 2013-08-06 21:55:56.956475885 +0000 @@ -31,7 +31,6 @@ void md5_crypt_gpu(crypt_md5_password *, 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, diff -urpN jumbo.orig/src/cuda_cryptmd5.h jumbo/src/cuda_cryptmd5.h --- jumbo.orig/src/cuda_cryptmd5.h 2013-08-05 23:59:46.159732736 +0000 +++ jumbo/src/cuda_cryptmd5.h 2013-08-06 21:39:12.926982806 +0000 @@ -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))){ diff -urpN jumbo.orig/src/cuda_phpass_fmt.c jumbo/src/cuda_phpass_fmt.c --- jumbo.orig/src/cuda_phpass_fmt.c 2013-08-05 23:59:46.159732736 +0000 +++ jumbo/src/cuda_phpass_fmt.c 2013-08-06 22:40:34.458744821 +0000 @@ -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, ¤tsalt, 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, ¤tsalt, 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, diff -urpN jumbo.orig/src/cuda_phpass.h jumbo/src/cuda_phpass.h --- jumbo.orig/src/cuda_phpass.h 2013-08-05 23:59:46.171732807 +0000 +++ jumbo/src/cuda_phpass.h 2013-08-06 00:27:32.924778770 +0000 @@ -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