diff -urpN john-1.7.8.orig//src/Makefile john-1.7.8-mscash2cuda-0/src/Makefile --- john-1.7.8.orig//src/Makefile 2011-05-04 18:52:31.000000000 +0000 +++ john-1.7.8-mscash2cuda-0/src/Makefile 2011-08-22 18:36:22.323678649 +0000 @@ -3,6 +3,7 @@ # Copyright (c) 1996-2011 by Solar Designer # +NVCC = nvcc CC = gcc AS = $(CC) LD = $(CC) @@ -19,6 +20,12 @@ OMPFLAGS = #OMPFLAGS = -fopenmp # Sun Studio with OpenMP (set the OMP_NUM_THREADS env var at runtime) #OMPFLAGS = -xopenmp +CUDAPATH=/usr/local/cuda/lib +CUDA64PATH=/usr/local/cuda/lib64 + +NVCC_FLAGS = -c -Xptxas -v -arch sm_10 +MSCASH_FLAGS = $(NVCC_FLAGS) --maxrregcount=60 + CFLAGS = -c -Wall -O2 -fomit-frame-pointer $(OMPFLAGS) ASFLAGS = -c $(OMPFLAGS) LDFLAGS = -s $(OMPFLAGS) @@ -41,6 +48,9 @@ JOHN_OBJS = \ unafs.o \ unique.o +CUDA_OBJS = \ + mscash2cuda_fmt.o mscash2.o + BENCH_DES_OBJS_ORIG = \ DES_fmt.o DES_std.o @@ -74,6 +84,8 @@ default: @echo "To build John the Ripper, type:" @echo " make clean SYSTEM" @echo "where SYSTEM can be one of the following:" + @echo "linux-x86-64-cuda Linux, x86-64 with SSE2 and CUDA (experimental)" + @echo "linux-x86-cuda Linux, x86 32-bit with SSE2 and CUDA (experimental)" @echo "linux-x86-64 Linux, x86-64 with SSE2 (best tested)" @echo "linux-x86-64-avx Linux, x86-64 with AVX (experimental)" @echo "linux-x86-64-xop Linux, x86-64 with AVX and XOP (experimental)" @@ -148,6 +160,20 @@ default: @echo "beos-x86-any BeOS, x86" @echo "generic Any other Unix-like system with gcc" +linux-x86-64-cuda: + $(LN) x86-64.h arch.h + $(MAKE) $(PROJ) \ + JOHN_OBJS="$(JOHN_OBJS) $(CUDA_OBJS) c3_fmt.o x86-64.o" \ + CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \ + LDFLAGS="$(LDFLAGS) -L$(CUDA64PATH) -lcudart -lcrypt" + +linux-x86-cuda: + $(LN) x86-sse.h arch.h + $(MAKE) $(PROJ) \ + JOHN_OBJS="$(JOHN_OBJS) $(CUDA_OBJS) c3_fmt.o x86.o x86-sse.o" \ + CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \ + LDFLAGS="$(LDFLAGS) -L$(CUDAPATH) -lcudart -lcrypt" + linux-x86-64: $(LN) x86-64.h arch.h $(MAKE) $(PROJ) \ @@ -811,6 +837,13 @@ generic.h: "$(BENCH_MD5_OBJS_DEPEND)" \ "$(BENCH_BF_OBJS_DEPEND)" +mscash2.o: mscash2.h cuda/mscash2.cu + cd cuda; nvcc $(MSCASH_FLAGS) mscash2.cu + +mscash2cuda_fmt.o: mscash2.o mscash2cuda_fmt.c mscash2.h + cp cuda/mscash2.o mscash2.o + $(CC) $(CFLAGS) mscash2cuda_fmt.c + bench: $(BENCH_OBJS) $(LD) $(BENCH_OBJS) $(LDFLAGS) -o bench @@ -890,7 +923,7 @@ depend: clean: $(RM) $(PROJ) $(PROJ_DOS) $(PROJ_WIN32) - $(RM) ../run/john.exe john-macosx-* *.o *.bak core + $(RM) ../run/john.exe john-macosx-* *.o *.bak *~ core cuda/*.o cuda/*~ $(RM) detect bench generic.h arch.h tmp.s $(CP) $(NULL) Makefile.dep diff -urpN john-1.7.8.orig//src/cuda/mscash2.cu john-1.7.8-mscash2cuda-0/src/cuda/mscash2.cu --- john-1.7.8.orig//src/cuda/mscash2.cu 1970-01-01 00:00:00.000000000 +0000 +++ john-1.7.8-mscash2cuda-0/src/cuda/mscash2.cu 2011-08-22 04:28:45.240396371 +0000 @@ -0,0 +1,408 @@ +/* +* This software is Copyright (c) 2011 Lukas Odzioba +* 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. +* Based on S3nf implementation http://openwall.info/wiki/john/MSCash2 +*/ + +#include +#include "../mscash2.h" + +static void HandleError(cudaError_t err, const char *file, int line) +{ + if (err != cudaSuccess) { + printf("%s in %s at line %d\n", cudaGetErrorString(err), file, + line); + exit(EXIT_FAILURE); + } +} + +#define HANDLE_ERROR(err) (HandleError(err,__FILE__,__LINE__)) + +extern "C" void mscash2_init(int gpuid); +extern "C" void mscash2_gpu(mscash2_password *, mscash2_hash *, mscash2_salt *); + +__constant__ mscash2_salt cuda_salt[1]; +__host__ void mscash2_init(int gpuid) +{ + int count; + HANDLE_ERROR(cudaGetDeviceCount(&count)); + if (gpuid < count) + cudaSetDevice(gpuid); + else { + printf("Invalid CUDA device id = %d\n", gpuid); + exit(1); + } +} + +__host__ void md4_crypt(uint32_t * buffer, uint32_t * hash) +{ + uint32_t a; + uint32_t b; + uint32_t c; + uint32_t d; + + a = 0xFFFFFFFF + buffer[0]; + a = (a << 3) | (a >> 29); + d = INIT_D + (INIT_C ^ (a & 0x77777777)) + buffer[1]; + d = (d << 7) | (d >> 25); + c = INIT_C + (INIT_B ^ (d & (a ^ INIT_B))) + buffer[2]; + c = (c << 11) | (c >> 21); + b = INIT_B + (a ^ (c & (d ^ a))) + buffer[3]; + b = (b << 19) | (b >> 13); + + a += (d ^ (b & (c ^ d))) + buffer[4]; + a = (a << 3) | (a >> 29); + d += (c ^ (a & (b ^ c))) + buffer[5]; + d = (d << 7) | (d >> 25); + c += (b ^ (d & (a ^ b))) + buffer[6]; + c = (c << 11) | (c >> 21); + b += (a ^ (c & (d ^ a))) + buffer[7]; + b = (b << 19) | (b >> 13); + + a += (d ^ (b & (c ^ d))) + buffer[8]; + a = (a << 3) | (a >> 29); + d += (c ^ (a & (b ^ c))) + buffer[9]; + d = (d << 7) | (d >> 25); + c += (b ^ (d & (a ^ b))) + buffer[10]; + c = (c << 11) | (c >> 21); + b += (a ^ (c & (d ^ a))) + buffer[11]; + b = (b << 19) | (b >> 13); + + a += (d ^ (b & (c ^ d))) + buffer[12]; + a = (a << 3) | (a >> 29); + d += (c ^ (a & (b ^ c))) + buffer[13]; + d = (d << 7) | (d >> 25); + c += (b ^ (d & (a ^ b))) + buffer[14]; + c = (c << 11) | (c >> 21); + b += (a ^ (c & (d ^ a))) + buffer[15]; + b = (b << 19) | (b >> 13); + + a += ((b & (c | d)) | (c & d)) + buffer[0] + SQRT_2; + a = (a << 3) | (a >> 29); + d += ((a & (b | c)) | (b & c)) + buffer[4] + SQRT_2; + d = (d << 5) | (d >> 27); + c += ((d & (a | b)) | (a & b)) + buffer[8] + SQRT_2; + c = (c << 9) | (c >> 23); + b += ((c & (d | a)) | (d & a)) + buffer[12] + SQRT_2; + b = (b << 13) | (b >> 19); + + a += ((b & (c | d)) | (c & d)) + buffer[1] + SQRT_2; + a = (a << 3) | (a >> 29); + d += ((a & (b | c)) | (b & c)) + buffer[5] + SQRT_2; + d = (d << 5) | (d >> 27); + c += ((d & (a | b)) | (a & b)) + buffer[9] + SQRT_2; + c = (c << 9) | (c >> 23); + b += ((c & (d | a)) | (d & a)) + buffer[13] + SQRT_2; + b = (b << 13) | (b >> 19); + + a += ((b & (c | d)) | (c & d)) + buffer[2] + SQRT_2; + a = (a << 3) | (a >> 29); + d += ((a & (b | c)) | (b & c)) + buffer[6] + SQRT_2; + d = (d << 5) | (d >> 27); + c += ((d & (a | b)) | (a & b)) + buffer[10] + SQRT_2; + c = (c << 9) | (c >> 23); + b += ((c & (d | a)) | (d & a)) + buffer[14] + SQRT_2; + b = (b << 13) | (b >> 19); + + a += ((b & (c | d)) | (c & d)) + buffer[3] + SQRT_2; + a = (a << 3) | (a >> 29); + d += ((a & (b | c)) | (b & c)) + buffer[7] + SQRT_2; + d = (d << 5) | (d >> 27); + c += ((d & (a | b)) | (a & b)) + buffer[11] + SQRT_2; + c = (c << 9) | (c >> 23); + b += ((c & (d | a)) | (d & a)) + buffer[15] + SQRT_2; + b = (b << 13) | (b >> 19); + + a += (d ^ c ^ b) + buffer[0] + SQRT_3; + a = (a << 3) | (a >> 29); + d += (c ^ b ^ a) + buffer[8] + SQRT_3; + d = (d << 9) | (d >> 23); + c += (b ^ a ^ d) + buffer[4] + SQRT_3; + c = (c << 11) | (c >> 21); + b += (a ^ d ^ c) + buffer[12] + SQRT_3; + b = (b << 15) | (b >> 17); + + a += (d ^ c ^ b) + buffer[2] + SQRT_3; + a = (a << 3) | (a >> 29); + d += (c ^ b ^ a) + buffer[10] + SQRT_3; + d = (d << 9) | (d >> 23); + c += (b ^ a ^ d) + buffer[6] + SQRT_3; + c = (c << 11) | (c >> 21); + b += (a ^ d ^ c) + buffer[14] + SQRT_3; + b = (b << 15) | (b >> 17); + + a += (d ^ c ^ b) + buffer[1] + SQRT_3; + a = (a << 3) | (a >> 29); + d += (c ^ b ^ a) + buffer[9] + SQRT_3; + d = (d << 9) | (d >> 23); + c += (b ^ a ^ d) + buffer[5] + SQRT_3; + c = (c << 11) | (c >> 21); + b += (a ^ d ^ c) + buffer[13] + SQRT_3; + b = (b << 15) | (b >> 17); + + a += (d ^ c ^ b) + buffer[3] + SQRT_3; + a = (a << 3) | (a >> 29); + + d += (c ^ b ^ a) + buffer[11] + SQRT_3; + d = (d << 9) | (d >> 23); + c += (b ^ a ^ d) + buffer[7] + SQRT_3; + c = (c << 11) | (c >> 21); + b += (a ^ d ^ c) + buffer[15] + SQRT_3; + b = (b << 15) | (b >> 17); + + hash[0] = a + INIT_A; + hash[1] = b + INIT_B; + hash[2] = c + INIT_C; + hash[3] = d + INIT_D; +} + +__device__ __host__ void preproc(const uint8_t * key, uint32_t keylen, + uint32_t * state, uint8_t var) +{ + int i; + uint32_t W[16], temp; + uint8_t ipad[64]; + + for (i = 0; i < 64; i++) + ipad[i] = var; + + for (i = 0; i < keylen; i++) + ipad[i] = ipad[i] ^ key[i]; + +#pragma unroll 16 + for (i = 0; i < 16; i++) + GET_WORD_32_BE(W[i], ipad, i * 4); + + uint32_t A = INIT_A; + uint32_t B = INIT_B; + uint32_t C = INIT_C; + uint32_t D = INIT_D; + uint32_t E = INIT_E; + + SHA1(A, B, C, D, E, W); + + state[0] = A + INIT_A; + state[1] = B + INIT_B; + state[2] = C + INIT_C; + state[3] = D + INIT_D; + state[4] = E + INIT_E; + +} + +__device__ void hmac_sha1(const uint8_t * key, uint32_t keylen, + const uint8_t * input, uint32_t inputlen, uint8_t * output, + uint32_t * ipad_state, uint32_t * opad_state) +{ + int i; + uint32_t temp, W[16]; + uint32_t A, B, C, D, E; + uint32_t state_A,state_B,state_C,state_D,state_E; + uint8_t buf[64]; + uint32_t *src=(uint32_t*)buf; + i=64/4; + while(i--) + *src++=0; + + memcpy(buf, input, inputlen); + buf[inputlen] = 0x80; + PUT_WORD_32_BE((64 + inputlen) << 3, buf, 60); + + A = ipad_state[0]; + B = ipad_state[1]; + C = ipad_state[2]; + D = ipad_state[3]; + E = ipad_state[4]; + + state_A=A; + state_B=B; + state_C=C; + state_D=D; + state_E=E; + + for (i = 0; i < 16; i++) + GET_WORD_32_BE(W[i], buf, i * 4); + + SHA1(A, B, C, D, E, W); + + A += state_A; + B += state_B; + C += state_C; + D += state_D; + E += state_E; + + PUT_WORD_32_BE(A, buf, 0); + PUT_WORD_32_BE(B, buf, 4); + PUT_WORD_32_BE(C, buf, 8); + PUT_WORD_32_BE(D, buf, 12); + PUT_WORD_32_BE(E, buf, 16); + + buf[20] = 0x80; + PUT_WORD_32_BE(0x2A0, buf, 60); + + A = opad_state[0]; + B = opad_state[1]; + C = opad_state[2]; + D = opad_state[3]; + E = opad_state[4]; + + state_A=A; + state_B=B; + state_C=C; + state_D=D; + state_E=E; + + for (i = 0; i < 16; i++) + GET_WORD_32_BE(W[i], buf, i * 4); + + SHA1(A, B, C, D, E, W); + + A += state_A; + B += state_B; + C += state_C; + D += state_D; + E += state_E; + + PUT_WORD_32_BE(A, output, 0); + PUT_WORD_32_BE(B, output, 4); + PUT_WORD_32_BE(C, output, 8); + PUT_WORD_32_BE(D, output, 12); + PUT_WORD_32_BE(E, output, 16); +} + +__device__ void pbkdf2(const uint8_t * pass, const uint8_t * salt, + int saltlen, uint8_t * out) +{ + uint8_t temp[SHA1_DIGEST_LENGTH]; + uint8_t buf[48]; + uint32_t ipad_state[5]; + uint32_t opad_state[5]; + int i, j; + uint8_t tmp_out[16]; + + i=48/4; + uint32_t *src=(uint32_t*)buf; + while(i--) + *src++=0; + + memcpy(buf, salt, saltlen); + buf[saltlen + 3] = 0x01; + + preproc(pass, 16, ipad_state, 0x36); + preproc(pass, 16, opad_state, 0x5c); + + hmac_sha1(pass, 16, buf, saltlen + 4, temp, ipad_state, opad_state); + + memcpy(tmp_out, temp, 20); + + for (i = 1; i < ITERATIONS; i++) { + hmac_sha1(pass, 16, temp, SHA1_DIGEST_LENGTH, temp, ipad_state, + opad_state); + +#pragma unroll 16 + for (j = 0; j < 16; j++) + tmp_out[j] ^= temp[j]; + } + memcpy(out, tmp_out, 20); +} + + +__global__ void pbkdf2_kernel(mscash2_password * inbuffer, + mscash2_hash * outbuffer) +{ + uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x; + + uint32_t username_len = (uint32_t) cuda_salt[0].length; + + pbkdf2((uint8_t *) inbuffer[idx].dcc_hash, + cuda_salt[0].unicode_salt, username_len << 1, + (uint8_t *) outbuffer[idx].v); + +} + +__host__ void mscash_cpu(mscash2_password * inbuffer, mscash2_hash * outbuffer, + mscash2_salt * host_salt) + { + + int i,idx = 0; + uint32_t buffer[16]; + uint32_t nt_hash[16]; + uint8_t salt[64]; + memset(salt,0,64); + uint8_t *username = host_salt->salt; + uint32_t username_len = (uint32_t) host_salt->length; + + + for (i = 0; i < (username_len >> 1) + 1; i++) + ((uint32_t *) salt)[i] = + username[2 * i] | (username[2 * i + 1] << 16); + memcpy(host_salt->unicode_salt, salt, 64); + + + + for (idx = 0; idx < KEYS_PER_CRYPT; idx++) { + + uint8_t *password = inbuffer[idx].v; + uint32_t password_len = inbuffer[idx].length; + memset(nt_hash, 0, 64); + memset(buffer, 0, 64); + + for (i = 0; i < password_len >> 1; i++) + buffer[i] = + password[2 * i] | (password[2 * i + 1] << 16); + + if (password_len % 2 == 1) + buffer[i] = password[password_len - 1] | 0x800000; + else + buffer[i] = 0x80; + + buffer[14] = password_len << 4; + + md4_crypt(buffer, nt_hash); + + memcpy((uint8_t *) nt_hash + 16, salt, username_len << 1); + + i = username_len + 8; + + if (username_len % 2 == 1) + nt_hash[i >> 1] = + username[username_len - 1] | 0x800000; + else + nt_hash[i >> 1] = 0x80; + + nt_hash[14] = i << 4; + + md4_crypt(nt_hash, inbuffer[idx].dcc_hash); + + } + + + } +__host__ void mscash2_gpu(mscash2_password * inbuffer, mscash2_hash * outbuffer, + mscash2_salt * host_salt) +{ + + mscash_cpu(inbuffer,outbuffer,host_salt); + mscash2_password *cuda_inbuffer; + mscash2_hash *cuda_outbuffer; + size_t insize = sizeof(mscash2_password) * KEYS_PER_CRYPT; + size_t outsize = sizeof(mscash2_hash) * KEYS_PER_CRYPT; + + HANDLE_ERROR(cudaMemcpyToSymbol(cuda_salt, host_salt, + sizeof(mscash2_salt))); + + HANDLE_ERROR(cudaMalloc(&cuda_inbuffer, insize)); + HANDLE_ERROR(cudaMalloc(&cuda_outbuffer, outsize)); + + HANDLE_ERROR(cudaMemcpy(cuda_inbuffer, inbuffer, insize, + cudaMemcpyHostToDevice)); + + pbkdf2_kernel <<< BLOCKS, THREADS >>> (cuda_inbuffer, cuda_outbuffer); + + HANDLE_ERROR(cudaMemcpy(outbuffer, cuda_outbuffer, outsize, + cudaMemcpyDeviceToHost)); + + HANDLE_ERROR(cudaFree(cuda_inbuffer)); + HANDLE_ERROR(cudaFree(cuda_outbuffer)); + +} diff -urpN john-1.7.8.orig//src/john.c john-1.7.8-mscash2cuda-0/src/john.c --- john-1.7.8.orig//src/john.c 2011-02-27 12:31:36.000000000 +0000 +++ john-1.7.8-mscash2cuda-0/src/john.c 2011-08-22 02:09:22.361521212 +0000 @@ -39,6 +39,7 @@ extern int CPU_detect(void); extern struct fmt_main fmt_DES, fmt_BSDI, fmt_MD5, fmt_BF; extern struct fmt_main fmt_AFS, fmt_LM; +extern struct fmt_main fmt_MSCASH2CUDA; #ifdef HAVE_CRYPT extern struct fmt_main fmt_crypt; #endif @@ -71,6 +72,7 @@ static void john_register_all(void) john_register_one(&fmt_BF); john_register_one(&fmt_AFS); john_register_one(&fmt_LM); + john_register_one(&fmt_MSCASH2CUDA); #ifdef HAVE_CRYPT john_register_one(&fmt_crypt); #endif diff -urpN john-1.7.8.orig//src/mscash2.h john-1.7.8-mscash2cuda-0/src/mscash2.h --- john-1.7.8.orig//src/mscash2.h 1970-01-01 00:00:00.000000000 +0000 +++ john-1.7.8-mscash2cuda-0/src/mscash2.h 2011-08-22 18:36:11.591678614 +0000 @@ -0,0 +1,202 @@ +/* +* This software is Copyright (c) 2011 Lukas Odzioba +* 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. +* Based on S3nf implementation http://openwall.info/wiki/john/MSCash2 +*/ +#ifndef _MSCASH2_H +#define _MSCASH2_H + +#define uint8_t unsigned char +#define uint16_t unsigned short +#define uint32_t unsigned int + +#define THREADS 128 +#define BLOCKS 14 +#define KEYS_PER_CRYPT (THREADS)*(BLOCKS) + +#define BINARY_SIZE 16 +#define PLAINTEXT_LENGTH 15 +#define SALT_SIZE sizeof(mscash2_salt) + +#define MIN_KEYS_PER_CRYPT KEYS_PER_CRYPT +#define MAX_KEYS_PER_CRYPT KEYS_PER_CRYPT + +#define MAX(x,y) ((x) > (y) ? (x) : (y)) +#define MIN(x,y) ((x) < (y) ? (x) : (y)) +# define SWAP(n) \ + (((n) << 24) | (((n) & 0xff00) << 8) | (((n) >> 8) & 0xff00) | ((n) >> 24)) + +#define ITERATIONS 10240 + +#define INIT_A 0x67452301 +#define INIT_B 0xefcdab89 +#define INIT_C 0x98badcfe +#define INIT_D 0x10325476 +#define INIT_E 0xc3d2e1f0 + +#define SQRT_2 0x5a827999 +#define SQRT_3 0x6ed9eba1 + +#define SHA1_DIGEST_LENGTH 20 + +#define K1 0x5a827999 +#define K2 0x6ed9eba1 +#define K3 0x8f1bbcdc +#define K4 0xca62c1d6 + +#define F1(x,y,z) (z ^ (x & (y ^ z))) +#define F2(x,y,z) (x ^ y ^ z) +#define F3(x,y,z) ((x & y) | (z & (x | y))) +#define F4(x,y,z) (x ^ y ^ z) + +#ifndef GET_WORD_32_BE +#define GET_WORD_32_BE(n,b,i) \ +{ \ + (n) = ( (unsigned long) (b)[(i) ] << 24 ) \ + | ( (unsigned long) (b)[(i) + 1] << 16 ) \ + | ( (unsigned long) (b)[(i) + 2] << 8 ) \ + | ( (unsigned long) (b)[(i) + 3] ); \ +} +#endif + +#ifndef PUT_WORD_32_BE +#define PUT_WORD_32_BE(n,b,i) \ +{ \ + (b)[(i) ] = (unsigned char) ( (n) >> 24 ); \ + (b)[(i) + 1] = (unsigned char) ( (n) >> 16 ); \ + (b)[(i) + 2] = (unsigned char) ( (n) >> 8 ); \ + (b)[(i) + 3] = (unsigned char) ( (n) ); \ +} +#endif + +#define S(x,n) ((x << n) | ((x & 0xFFFFFFFF) >> (32 - n))) + +#define R(t) \ +( \ + temp = W[(t - 3) & 0x0F] ^ W[(t - 8) & 0x0F] ^ \ + W[(t - 14) & 0x0F] ^ W[ t & 0x0F], \ + ( W[t & 0x0F] = S(temp,1) ) \ +) + +#define P1(a,b,c,d,e,x) \ +{ \ + e += S(a,5) + F1(b,c,d) + K1 + x; b = S(b,30); \ +} + +#define P2(a,b,c,d,e,x) \ +{ \ + e += S(a,5) + F2(b,c,d) + K2 + x; b = S(b,30); \ +} + +#define P3(a,b,c,d,e,x) \ +{ \ + e += S(a,5) + F3(b,c,d) + K3 + x; b = S(b,30); \ +} + +#define P4(a,b,c,d,e,x) \ +{ \ + e += S(a,5) + F4(b,c,d) + K4 + x; b = S(b,30); \ +} + +#define SHA1(A,B,C,D,E,W) \ + P1(A, B, C, D, E, W[0] );\ + P1(E, A, B, C, D, W[1] );\ + P1(D, E, A, B, C, W[2] );\ + P1(C, D, E, A, B, W[3] );\ + P1(B, C, D, E, A, W[4] );\ + P1(A, B, C, D, E, W[5] );\ + P1(E, A, B, C, D, W[6] );\ + P1(D, E, A, B, C, W[7] );\ + P1(C, D, E, A, B, W[8] );\ + P1(B, C, D, E, A, W[9] );\ + P1(A, B, C, D, E, W[10]);\ + P1(E, A, B, C, D, W[11]);\ + P1(D, E, A, B, C, W[12]);\ + P1(C, D, E, A, B, W[13]);\ + P1(B, C, D, E, A, W[14]);\ + P1(A, B, C, D, E, W[15]);\ + P1(E, A, B, C, D, R(16));\ + P1(D, E, A, B, C, R(17));\ + P1(C, D, E, A, B, R(18));\ + P1(B, C, D, E, A, R(19));\ + P2(A, B, C, D, E, R(20));\ + P2(E, A, B, C, D, R(21));\ + P2(D, E, A, B, C, R(22));\ + P2(C, D, E, A, B, R(23));\ + P2(B, C, D, E, A, R(24));\ + P2(A, B, C, D, E, R(25));\ + P2(E, A, B, C, D, R(26));\ + P2(D, E, A, B, C, R(27));\ + P2(C, D, E, A, B, R(28));\ + P2(B, C, D, E, A, R(29));\ + P2(A, B, C, D, E, R(30));\ + P2(E, A, B, C, D, R(31));\ + P2(D, E, A, B, C, R(32));\ + P2(C, D, E, A, B, R(33));\ + P2(B, C, D, E, A, R(34));\ + P2(A, B, C, D, E, R(35));\ + P2(E, A, B, C, D, R(36));\ + P2(D, E, A, B, C, R(37));\ + P2(C, D, E, A, B, R(38));\ + P2(B, C, D, E, A, R(39));\ + P3(A, B, C, D, E, R(40));\ + P3(E, A, B, C, D, R(41));\ + P3(D, E, A, B, C, R(42));\ + P3(C, D, E, A, B, R(43));\ + P3(B, C, D, E, A, R(44));\ + P3(A, B, C, D, E, R(45));\ + P3(E, A, B, C, D, R(46));\ + P3(D, E, A, B, C, R(47));\ + P3(C, D, E, A, B, R(48));\ + P3(B, C, D, E, A, R(49));\ + P3(A, B, C, D, E, R(50));\ + P3(E, A, B, C, D, R(51));\ + P3(D, E, A, B, C, R(52));\ + P3(C, D, E, A, B, R(53));\ + P3(B, C, D, E, A, R(54));\ + P3(A, B, C, D, E, R(55));\ + P3(E, A, B, C, D, R(56));\ + P3(D, E, A, B, C, R(57));\ + P3(C, D, E, A, B, R(58));\ + P3(B, C, D, E, A, R(59));\ + P4(A, B, C, D, E, R(60));\ + P4(E, A, B, C, D, R(61));\ + P4(D, E, A, B, C, R(62));\ + P4(C, D, E, A, B, R(63));\ + P4(B, C, D, E, A, R(64));\ + P4(A, B, C, D, E, R(65));\ + P4(E, A, B, C, D, R(66));\ + P4(D, E, A, B, C, R(67));\ + P4(C, D, E, A, B, R(68));\ + P4(B, C, D, E, A, R(69));\ + P4(A, B, C, D, E, R(70));\ + P4(E, A, B, C, D, R(71));\ + P4(D, E, A, B, C, R(72));\ + P4(C, D, E, A, B, R(73));\ + P4(B, C, D, E, A, R(74));\ + P4(A, B, C, D, E, R(75));\ + P4(E, A, B, C, D, R(76));\ + P4(D, E, A, B, C, R(77));\ + P4(C, D, E, A, B, R(78));\ + P4(B, C, D, E, A, R(79)); + +static const char mscash2_prefix[] = "$DCC2$"; + +typedef struct { + uint8_t length; + uint8_t v[15]; + uint32_t dcc_hash[4]; +} mscash2_password; + +typedef struct { + uint32_t v[8]; +} mscash2_hash; + +typedef struct { + uint8_t length; + uint8_t salt[15]; + uint8_t unicode_salt[64]; +} mscash2_salt; + +#endif diff -urpN john-1.7.8.orig//src/mscash2cuda_fmt.c john-1.7.8-mscash2cuda-0/src/mscash2cuda_fmt.c --- john-1.7.8.orig//src/mscash2cuda_fmt.c 1970-01-01 00:00:00.000000000 +0000 +++ john-1.7.8-mscash2cuda-0/src/mscash2cuda_fmt.c 2011-09-09 06:55:01.174592599 +0000 @@ -0,0 +1,244 @@ +/* +* This software is Copyright (c) 2011 Lukas Odzioba +* 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. +* Based on S3nf implementation http://openwall.info/wiki/john/MSCash2 +*/ +#include +#include "arch.h" +#include "formats.h" +#include "common.h" +#include "misc.h" +#include "mscash2.h" + +#define FORMAT_LABEL "mscash2cuda" +#define FORMAT_NAME "MSCASH2CUDA" +#define ALGORITHM_NAME "GPU" + +#define BENCHMARK_COMMENT "" +#define BENCHMARK_LENGTH -1 + +//#define _MSCASH2_DEBUG + +static mscash2_password inbuffer[MAX_KEYS_PER_CRYPT]; +static mscash2_hash outbuffer[MAX_KEYS_PER_CRYPT]; +static mscash2_salt currentsalt; + +static struct fmt_tests tests[] = { + {"$DCC2$test#a86012faf7d88d1fc037a69764a92cac", "password"}, + //{"$DCC2$administrator#a150f71752b5d605ef0b2a1e98945611","a"}, + //{"$DCC2$administrator#c14eb8279e4233ec14e9d393637b65e2","ab"}, + //{"$DCC2$administrator#8ce9c0279b4e6f226f52d559f9c2c5f3","abc"}, + //{"$DCC2$administrator#2fc788d09fad7e26a92d12356fa44bdf","abcd"}, + //{"$DCC2$administrator#6aa19842ffea11f0f0c89f8ca8d245bd","abcde"}, + {NULL} +}; + +extern void mscash2_init(int gpuid); +extern void mscash2_gpu(mscash2_password *, mscash2_hash *, mscash2_salt *); + +static void init(void) +{ + mscash2_init(0); +} + +static int valid(char *ciphertext) +{ + if (strncmp(ciphertext, mscash2_prefix, strlen(mscash2_prefix)) != 0) + return 0; + char *hash = strrchr(ciphertext, '#') + 1; + int hashlength = 0; + if (hash == NULL) + return 0; + while (hash < ciphertext + strlen(ciphertext)) { + if (atoi16[ARCH_INDEX(*hash++)] == 0x7f) + return 0; + hashlength++; + } + if (hashlength != 32) + return 0; + return 1; +} + +static void *binary(char *ciphertext) +{ + static uint32_t binary[4]; + char *hash = strrchr(ciphertext, '#') + 1; + if (hash == NULL) + return binary; + int i; + for (i = 0; i < 4; i++) { + sscanf(hash + (8 * i), "%08x", &binary[i]); + binary[i] = SWAP(binary[i]); + } + return binary; + +} + +static void *salt(char *ciphertext) +{ + static mscash2_salt salt; + char *pos = ciphertext + strlen(mscash2_prefix); + int length = 0; + while (*pos != '#') + salt.salt[length++] = *pos++; + salt.length = length; + return &salt; +} + +static void set_salt(void *salt) +{ + memcpy(¤tsalt, salt, sizeof(mscash2_salt)); +} + +static void set_key(char *key, int index) +{ + uint8_t length = strlen(key); + inbuffer[index].length = length; + memcpy(inbuffer[index].v, key, length); +} + +static char *get_key(int index) +{ + static char ret[PLAINTEXT_LENGTH + 1]; + uint8_t length = inbuffer[index].length; + memcpy(ret, inbuffer[index].v, length); + ret[length] = '\0'; + return ret; +} + +static void crypt_all(int count) +{ + mscash2_gpu(inbuffer, outbuffer, ¤tsalt); +} + +static int binary_hash_0(void *binary) +{ +#ifdef _MSCASH2_DEBUG + puts("binary"); + uint32_t i, *b = binary; + for (i = 0; i < 4; i++) + printf("%08x ", b[i]); + puts(""); +#endif + return (((uint32_t *) binary)[0] & 0xf); +} + +static int binary_hash_1(void *binary) +{ + return ((uint32_t *) binary)[0] & 0xff; +} + +static int binary_hash_2(void *binary) +{ + return ((uint32_t *) binary)[0] & 0xfff; +} + +static int binary_hash_3(void *binary) +{ + return ((uint32_t *) binary)[0] & 0xffff; +} + +static int binary_hash_4(void *binary) +{ + return ((uint32_t *) binary)[0] & 0xfffff; +} + + +static int get_hash_0(int index) +{ +#ifdef _MSCASH2_DEBUG + int i; + puts("get_hash"); + for (i = 0; i < 4; i++) + printf("%08x ", outbuffer[index].v[i]); + puts(""); +#endif + return outbuffer[index].v[0] & 0xf; +} + +static int get_hash_1(int index) +{ + return outbuffer[index].v[0] & 0xff; +} + +static int get_hash_2(int index) +{ + return outbuffer[index].v[0] & 0xfff; +} + +static int get_hash_3(int index) +{ + return outbuffer[index].v[0] & 0xffff; +} + +static int get_hash_4(int index) +{ + return outbuffer[index].v[0] & 0xfffff; +} + +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[i].v[0]) + return 1; + return 0; +} + +static int cmp_one(void *binary, int index) +{ + uint32_t i, *b = (uint32_t *) binary; + for (i = 0; i < 4; i++) + if (b[i] != outbuffer[index].v[i]) + return 0; + return 1; +} + +static int cmp_exact(char *source, int count) +{ + return 1; +} + +struct fmt_main fmt_MSCASH2CUDA = { + { + FORMAT_LABEL, + FORMAT_NAME, + ALGORITHM_NAME, + BENCHMARK_COMMENT, + BENCHMARK_LENGTH, + PLAINTEXT_LENGTH, + BINARY_SIZE, + SALT_SIZE, + MIN_KEYS_PER_CRYPT, + MAX_KEYS_PER_CRYPT, + FMT_CASE | FMT_8_BIT, + tests}, + { + init, + valid, + fmt_default_split, + binary, + salt, + { + binary_hash_0, + binary_hash_1, + binary_hash_2, + binary_hash_3, + binary_hash_4}, + fmt_default_salt_hash, + set_salt, + set_key, + get_key, + fmt_default_clear_keys, + crypt_all, + { + get_hash_0, + get_hash_1, + get_hash_2, + get_hash_3, + get_hash_4}, + cmp_all, + cmp_one, + cmp_exact} +}; diff -urpN john-1.7.8.orig//src/options.c john-1.7.8-mscash2cuda-0/src/options.c --- john-1.7.8.orig//src/options.c 2011-06-22 13:01:40.000000000 +0000 +++ john-1.7.8-mscash2cuda-0/src/options.c 2011-08-22 02:09:08.630270899 +0000 @@ -100,7 +100,7 @@ static struct opt_entry opt_list[] = { "--salts=[-]COUNT load salts with[out] at least COUNT passwords " \ "only\n" \ "--format=NAME force hash type NAME: " \ - "DES/BSDI/MD5/BF/AFS/LM" MAYBE_CRYPT "\n" \ + "DES/BSDI/MD5/BF/AFS/LM/MSCASH2CUDA" MAYBE_CRYPT "\n" \ "--save-memory=LEVEL enable memory saving, at LEVEL 1..3\n" void opt_init(char *name, int argc, char **argv)