diff -urpN john-1.7.8.orig//src/Makefile john-1.7.8-mscash-1//src/Makefile --- john-1.7.8.orig//src/Makefile 2011-05-04 18:52:31.000000000 +0000 +++ john-1.7.8-mscash-1//src/Makefile 2011-08-16 21:02:41.728774905 +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) + 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 = \ + mscashcuda_fmt.o mscash.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)" +mscash.o: mscash.h cuda/mscash.cu + cd cuda; nvcc $(MSCASH_FLAGS) mscash.cu + +mscashcuda_fmt.o: mscash.o mscashcuda_fmt.c mscash.h + cp cuda/mscash.o mscash.o + $(CC) $(CFLAGS) mscashcuda_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 $(RM) detect bench generic.h arch.h tmp.s $(CP) $(NULL) Makefile.dep diff -urpN john-1.7.8.orig//src/cuda/mscash.cu john-1.7.8-mscash-1//src/cuda/mscash.cu --- john-1.7.8.orig//src/cuda/mscash.cu 1970-01-01 00:00:00.000000000 +0000 +++ john-1.7.8-mscash-1//src/cuda/mscash.cu 2011-08-16 20:56:32.400743360 +0000 @@ -0,0 +1,248 @@ +/* +* 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 Alain Espinosa implementation http://openwall.info/wiki/john/MSCash +*/ + +#include +#include "../mscash.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 mscash_init(int gpuid); +extern "C" void mscash_gpu(mscash_password *, mscash_hash *, mscash_salt *); + +__constant__ mscash_salt cuda_salt[1]; +__host__ void mscash_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); + } +} + + +__device__ void md4_crypt(uint32_t * output, uint32_t * nt_buffer) +{ + unsigned int a = INIT_A; + unsigned int b = INIT_B; + unsigned int c = INIT_C; + unsigned int d = INIT_D; + + /* Round 1 */ + a += (d ^ (b & (c ^ d))) + nt_buffer[0]; + a = (a << 3) | (a >> 29); + d += (c ^ (a & (b ^ c))) + nt_buffer[1]; + d = (d << 7) | (d >> 25); + c += (b ^ (d & (a ^ b))) + nt_buffer[2]; + c = (c << 11) | (c >> 21); + b += (a ^ (c & (d ^ a))) + nt_buffer[3]; + b = (b << 19) | (b >> 13); + + a += (d ^ (b & (c ^ d))) + nt_buffer[4]; + a = (a << 3) | (a >> 29); + d += (c ^ (a & (b ^ c))) + nt_buffer[5]; + d = (d << 7) | (d >> 25); + c += (b ^ (d & (a ^ b))) + nt_buffer[6]; + c = (c << 11) | (c >> 21); + b += (a ^ (c & (d ^ a))) + nt_buffer[7]; + b = (b << 19) | (b >> 13); + + a += (d ^ (b & (c ^ d))) + nt_buffer[8]; + a = (a << 3) | (a >> 29); + d += (c ^ (a & (b ^ c))) + nt_buffer[9]; + d = (d << 7) | (d >> 25); + c += (b ^ (d & (a ^ b))) + nt_buffer[10]; + c = (c << 11) | (c >> 21); + b += (a ^ (c & (d ^ a))) + nt_buffer[11]; + b = (b << 19) | (b >> 13); + + a += (d ^ (b & (c ^ d))) + nt_buffer[12]; + a = (a << 3) | (a >> 29); + d += (c ^ (a & (b ^ c))) + nt_buffer[13]; + d = (d << 7) | (d >> 25); + c += (b ^ (d & (a ^ b))) + nt_buffer[14]; + c = (c << 11) | (c >> 21); + b += (a ^ (c & (d ^ a))) + nt_buffer[15]; + b = (b << 19) | (b >> 13); + + /* Round 2 */ + a += ((b & (c | d)) | (c & d)) + nt_buffer[0] + SQRT_2; + a = (a << 3) | (a >> 29); + d += ((a & (b | c)) | (b & c)) + nt_buffer[4] + SQRT_2; + d = (d << 5) | (d >> 27); + c += ((d & (a | b)) | (a & b)) + nt_buffer[8] + SQRT_2; + c = (c << 9) | (c >> 23); + b += ((c & (d | a)) | (d & a)) + nt_buffer[12] + SQRT_2; + b = (b << 13) | (b >> 19); + + a += ((b & (c | d)) | (c & d)) + nt_buffer[1] + SQRT_2; + a = (a << 3) | (a >> 29); + d += ((a & (b | c)) | (b & c)) + nt_buffer[5] + SQRT_2; + d = (d << 5) | (d >> 27); + c += ((d & (a | b)) | (a & b)) + nt_buffer[9] + SQRT_2; + c = (c << 9) | (c >> 23); + b += ((c & (d | a)) | (d & a)) + nt_buffer[13] + SQRT_2; + b = (b << 13) | (b >> 19); + + a += ((b & (c | d)) | (c & d)) + nt_buffer[2] + SQRT_2; + a = (a << 3) | (a >> 29); + d += ((a & (b | c)) | (b & c)) + nt_buffer[6] + SQRT_2; + d = (d << 5) | (d >> 27); + c += ((d & (a | b)) | (a & b)) + nt_buffer[10] + SQRT_2; + c = (c << 9) | (c >> 23); + b += ((c & (d | a)) | (d & a)) + nt_buffer[14] + SQRT_2; + b = (b << 13) | (b >> 19); + + a += ((b & (c | d)) | (c & d)) + nt_buffer[3] + SQRT_2; + a = (a << 3) | (a >> 29); + d += ((a & (b | c)) | (b & c)) + nt_buffer[7] + SQRT_2; + d = (d << 5) | (d >> 27); + c += ((d & (a | b)) | (a & b)) + nt_buffer[11] + SQRT_2; + c = (c << 9) | (c >> 23); + b += ((c & (d | a)) | (d & a)) + nt_buffer[15] + SQRT_2; + b = (b << 13) | (b >> 19); + + /* Round 3 */ + a += (d ^ c ^ b) + nt_buffer[0] + SQRT_3; + a = (a << 3) | (a >> 29); + d += (c ^ b ^ a) + nt_buffer[8] + SQRT_3; + d = (d << 9) | (d >> 23); + c += (b ^ a ^ d) + nt_buffer[4] + SQRT_3; + c = (c << 11) | (c >> 21); + b += (a ^ d ^ c) + nt_buffer[12] + SQRT_3; + b = (b << 15) | (b >> 17); + + a += (d ^ c ^ b) + nt_buffer[2] + SQRT_3; + a = (a << 3) | (a >> 29); + d += (c ^ b ^ a) + nt_buffer[10] + SQRT_3; + d = (d << 9) | (d >> 23); + c += (b ^ a ^ d) + nt_buffer[6] + SQRT_3; + c = (c << 11) | (c >> 21); + b += (a ^ d ^ c) + nt_buffer[14] + SQRT_3; + b = (b << 15) | (b >> 17); + + a += (d ^ c ^ b) + nt_buffer[1] + SQRT_3; + a = (a << 3) | (a >> 29); + d += (c ^ b ^ a) + nt_buffer[9] + SQRT_3; + d = (d << 9) | (d >> 23); + c += (b ^ a ^ d) + nt_buffer[5] + SQRT_3; + c = (c << 11) | (c >> 21); + b += (a ^ d ^ c) + nt_buffer[13] + SQRT_3; + b = (b << 15) | (b >> 17); + + a += (d ^ c ^ b) + nt_buffer[3] + SQRT_3; + a = (a << 3) | (a >> 29); + d += (c ^ b ^ a) + nt_buffer[11] + SQRT_3; + d = (d << 9) | (d >> 23); + c += (b ^ a ^ d) + nt_buffer[7] + SQRT_3; + c = (c << 11) | (c >> 21); + b += (a ^ d ^ c) + nt_buffer[15] + SQRT_3; + b = (b << 15) | (b >> 17); + + output[0] = a + INIT_A; + output[1] = b + INIT_B; + output[2] = c + INIT_C; + output[3] = d + INIT_D; +} + +__device__ void prepare_key(uint8_t * key, int length, uint32_t * nt_buffer) +{ + int i = 0; + for (i = 0; i < 16; i++) + nt_buffer[i] = 0; + for (i = 0; i < length / 2; i++) + nt_buffer[i] = key[2 * i] | (key[2 * i + 1] << 16); + if (length % 2 == 1) + nt_buffer[i] = key[length - 1] | 0x800000; + else + nt_buffer[i] = 0x80; + nt_buffer[14] = length << 4; +} + +__device__ void prepare_login(uint8_t * login, int length, + uint32_t * login_buffer) +{ + int i = 0; + for (i = 0; i < 12; i++) + login_buffer[i] = 0; + for (i = 0; i < length / 2; i++) + login_buffer[i] = login[2 * i] | (login[2 * i + 1] << 16); + if (length % 2 == 1) + login_buffer[i] = login[length - 1] | 0x800000; + else + login_buffer[i] = 0x80; + login_buffer[10] = (length << 4) + 128; +} + + +__global__ void mscash_kernel(mscash_password * inbuffer, + mscash_hash * outbuffer) +{ + uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x; + uint8_t *login = cuda_salt[0].salt; + uint8_t loginlength = cuda_salt[0].length; + + uint8_t *password = inbuffer[idx].v; + uint8_t passwordlength = inbuffer[idx].length; + + int i; + uint32_t nt_buffer[16]; + uint32_t login_buffer[12]; + uint32_t output[4]; + + prepare_key(password, passwordlength, nt_buffer); + md4_crypt(output, nt_buffer); + memcpy(nt_buffer, output, 4 * 4); + + prepare_login(login, loginlength, login_buffer); + memcpy(nt_buffer + 4, login_buffer, 12 * 4); + md4_crypt(output, nt_buffer); + + for (i = 0; i < 4; i++) + outbuffer[idx].v[i] = output[i]; + +} + +__host__ void mscash_gpu(mscash_password * inbuffer, mscash_hash * outbuffer, + mscash_salt * host_salt) +{ + + HANDLE_ERROR(cudaMemcpyToSymbol(cuda_salt, host_salt, + sizeof(mscash_salt))); + mscash_password *cuda_inbuffer; + mscash_hash *cuda_outbuffer; + + size_t insize = sizeof(mscash_password) * KEYS_PER_CRYPT; + size_t outsize = sizeof(mscash_hash) * KEYS_PER_CRYPT; + + HANDLE_ERROR(cudaMalloc(&cuda_inbuffer, insize)); + HANDLE_ERROR(cudaMalloc(&cuda_outbuffer, outsize)); + + HANDLE_ERROR(cudaMemcpy(cuda_inbuffer, inbuffer, insize, + cudaMemcpyHostToDevice)); + + mscash_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-mscash-1//src/john.c --- john-1.7.8.orig//src/john.c 2011-02-27 12:31:36.000000000 +0000 +++ john-1.7.8-mscash-1//src/john.c 2011-08-16 11:58:27.707771302 +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_MSCASHCUDA; #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_MSCASHCUDA); #ifdef HAVE_CRYPT john_register_one(&fmt_crypt); #endif diff -urpN john-1.7.8.orig//src/mscash.h john-1.7.8-mscash-1//src/mscash.h --- john-1.7.8.orig//src/mscash.h 1970-01-01 00:00:00.000000000 +0000 +++ john-1.7.8-mscash-1//src/mscash.h 2011-08-16 20:57:14.782646810 +0000 @@ -0,0 +1,56 @@ +/* +* 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 Alain Espinosa implementation http://openwall.info/wiki/john/MSCash +*/ +#ifndef _PBKDF2_H +#define _PBKDF2_H + +#define uint8_t unsigned char +#define uint16_t unsigned short +#define uint32_t unsigned int + +#define THREADS 512 +#define BLOCKS 21 +#define KEYS_PER_CRYPT (THREADS)*(BLOCKS) + +#define BINARY_SIZE 16 +#define PLAINTEXT_LENGTH 15 +#define SALT_SIZE sizeof(mscash_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 SHOW(x) (printf("%s = %08x\n",#x,(x))) +# define SWAP(n) \ + (((n) << 24) | (((n) & 0xff00) << 8) | (((n) >> 8) & 0xff00) | ((n) >> 24)) + +#define INIT_A 0x67452301 +#define INIT_B 0xefcdab89 +#define INIT_C 0x98badcfe +#define INIT_D 0x10325476 + +#define SQRT_2 0x5a827999 +#define SQRT_3 0x6ed9eba1 + +static const char mscash_prefix[] = "M$"; + +typedef struct { + uint8_t length; + uint8_t v[15]; +} mscash_password; + +typedef struct { + uint32_t v[8]; +} mscash_hash; + +typedef struct { + uint8_t length; + uint8_t salt[15]; +} mscash_salt; + +#endif diff -urpN john-1.7.8.orig//src/mscashcuda_fmt.c john-1.7.8-mscash-1//src/mscashcuda_fmt.c --- john-1.7.8.orig//src/mscashcuda_fmt.c 1970-01-01 00:00:00.000000000 +0000 +++ john-1.7.8-mscash-1//src/mscashcuda_fmt.c 2011-08-16 20:27:05.723720525 +0000 @@ -0,0 +1,217 @@ +/* +* 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 Alain Espinosa implementation http://openwall.info/wiki/john/MSCash +*/ +#include +#include "arch.h" +#include "formats.h" +#include "common.h" +#include "misc.h" +#include "mscash.h" + +#define FORMAT_LABEL "mscashcuda" +#define FORMAT_NAME "MSCASHCUDA" +#define ALGORITHM_NAME "" + +#define BENCHMARK_COMMENT "" +#define BENCHMARK_LENGTH -1 + +static mscash_password inbuffer[MAX_KEYS_PER_CRYPT]; +static mscash_hash outbuffer[MAX_KEYS_PER_CRYPT]; +static mscash_salt currentsalt; + +static struct fmt_tests tests[] = { + {"M$administrator#25fd08fa89795ed54207e6e8442a6ca0", "password"}, + {NULL} +}; + +extern void mscash_init(int gpuid); +extern void mscash_gpu(mscash_password *, mscash_hash *, mscash_salt *); + +static void init(void) +{ + mscash_init(0); +} + +static int valid(char *ciphertext) +{ +//valid prefix + if (strncmp(ciphertext, mscash_prefix, strlen(mscash_prefix)) != 0) + return 0; +//valid characters in hash + char *hash = strrchr(ciphertext, '#') + 1; + while (hash < ciphertext + strlen(ciphertext)) + if (atoi16[*hash++] == 0x7f) + return 0; + return 1; +} + +static void *binary(char *ciphertext) +{ + static uint32_t binary[4]; + char *hash = strrchr(ciphertext, '#') + 1; + 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 mscash_salt salt; + char *pos = ciphertext + strlen(mscash_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(mscash_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) +{ + mscash_gpu(inbuffer, outbuffer, ¤tsalt); +} + +static int binary_hash_0(void *binary) +{ + 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) +{ + 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_MSCASHCUDA = { + { + 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-mscash-1//src/options.c --- john-1.7.8.orig//src/options.c 2011-06-22 13:01:40.000000000 +0000 +++ john-1.7.8-mscash-1//src/options.c 2011-08-16 11:57:47.674896792 +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/MSCASHCUDA" MAYBE_CRYPT "\n" \ "--save-memory=LEVEL enable memory saving, at LEVEL 1..3\n" void opt_init(char *name, int argc, char **argv)