diff -urpN john-1.7.8.orig//doc/CRYPTMD5_CUDA john-1.7.8-cryptmd5_7//doc/CRYPTMD5_CUDA
--- john-1.7.8.orig//doc/CRYPTMD5_CUDA	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.8-cryptmd5_7//doc/CRYPTMD5_CUDA	2011-07-25 05:51:20.148819020 +0000
@@ -0,0 +1,75 @@
+  ,-*
+ (_)   Release Notes
+
+==== Version        ====
+    CRYPTMD5 CUDA ver. 3.0 patch for John The Ripper 1.7.8
+
+==== Release Date   ====
+    DD.MM.RRRR = 25.07.2011
+
+==== Authors        ====
+    Lukas Odzioba, lukas[dot]odzioba[at]gmail[dot]com
+
+==== Change Log     ====
+      ver 3:
+	Minor performance changes
+	Code cleanups
+	salt() bug fixed
+	$apr1$ support added
+      ver 2:
+	Shared memory utilization
+      ver 1:
+	Performance changes
+      ver 0:
+	First revision
+       
+==== Restrictions   ====
+    Patch was tested/optimized for following configuration:
+    i3-2100, 4GB 1333Mhz, Slackware 13.1, GTX 460 1GB
+    Patch is not setting up THREADS and BLOCKS automaticaly, you should experiment with those constants 
+    to get better performance on your card.
+    Passwords length should be less than 16.
+
+==== ToDo 	    ====
+    Speed improvements
+
+==== Known Bugs     ====
+
+==== Environment    ====
+    CRYPTMD5CUDA is supported/tested only on following builds:
+      linux-x86-sse2, linux-x86-mmx, linux-x86-64, linux-x86-any
+    If you have problems while building patched JtR:
+      Please check that paths to yours CUDA library are proper in Makefile lines:
+	24: CUDAPATH = /usr/local/cuda/lib
+	25: CUDA64PATH = /usr/local/cuda/lib64
+      Please check that yours PATH contains cuda lib,bin and include paths.
+      Please check that yours LD_LIBRARY_PATH contains cuda lib path.
+
+==== Instalation    ====
+    If you have an older GPU based on G80 or GT200 architecture change arch=sm_20 to sm_10 in the Makefile
+    Apply this patch on top of JtR 1.7.8
+    $ cd john-1.7.8
+    $ patch -p1 < ../john-1.7.8-cryptmd5cuda-3.diff
+    $ cd john-1.7.8/src
+    for 32-bit systems:
+    $ make clean && make linux-x86-sse2
+    for 64-bit systems:
+    $ make clean && make linux-x86-64
+
+==== Usage          ====
+    Benchmark:
+    $ john-1.7.8/run/john -test --format=cryptmd5cuda
+
+    Incremental attack:
+    $ john-1.7.8/run/john -i -format=cryptmd5cuda FILE
+
+==== Benchmarks     ==== 
+    ver 3:
+      GeForce GTX 460 1024MB		Raw:    391372 c/s real, 392157 c/s virtual
+    ver 2:
+      GeForce GTX 460 1024MB		Raw:    351260 c/s real, 355509 c/s virtual
+    ver 1:
+      GeForce GTX 460 1024MB		Raw:    143360 c/s real, 143360 c/s virtual
+    ver 0:
+      GeForce GTX 460 1024MB		Raw:    122328 c/s real, 122572 c/s virtual
+
diff -urpN john-1.7.8.orig//src/Makefile john-1.7.8-cryptmd5_7//src/Makefile
--- john-1.7.8.orig//src/Makefile	2011-05-04 18:52:31.000000000 +0000
+++ john-1.7.8-cryptmd5_7//src/Makefile	2011-07-25 05:55:54.674063776 +0000
@@ -4,6 +4,7 @@
 #
 
 CC = gcc
+NVCC = /usr/local/cuda/bin/nvcc
 AS = $(CC)
 LD = $(CC)
 CPP = $(CC)
@@ -19,6 +20,9 @@ OMPFLAGS =
 #OMPFLAGS = -fopenmp
 # Sun Studio with OpenMP (set the OMP_NUM_THREADS env var at runtime)
 #OMPFLAGS = -xopenmp
+NVCCFLAGS = -c -Xptxas -v -arch sm_20 
+CUDAPATH = /usr/local/cuda/lib
+CUDA64PATH = /usr/local/cuda/lib64
 CFLAGS = -c -Wall -O2 -fomit-frame-pointer $(OMPFLAGS)
 ASFLAGS = -c $(OMPFLAGS)
 LDFLAGS = -s $(OMPFLAGS)
@@ -41,6 +45,9 @@ JOHN_OBJS = \
 	unafs.o \
 	unique.o
 
+CUDA_OBJS = \
+	cryptmd5cuda_fmt.o cryptmd5.o 
+
 BENCH_DES_OBJS_ORIG = \
 	DES_fmt.o DES_std.o
 
@@ -151,9 +158,9 @@ default:
 linux-x86-64:
 	$(LN) x86-64.h arch.h
 	$(MAKE) $(PROJ) \
-		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86-64.o" \
+		JOHN_OBJS="$(JOHN_OBJS) $(CUDA_OBJS) c3_fmt.o x86-64.o" \
 		CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \
-		LDFLAGS="$(LDFLAGS) -lcrypt"
+		LDFLAGS="$(LDFLAGS) -L$(CUDA64PATH) -lcudart -lcrypt"
 
 linux-x86-64-avx:
 	$(LN) x86-64.h arch.h
@@ -188,23 +195,23 @@ linux-x86-64-32-mmx:
 linux-x86-sse2:
 	$(LN) x86-sse.h arch.h
 	$(MAKE) $(PROJ) \
-		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86.o x86-sse.o" \
+		JOHN_OBJS="$(JOHN_OBJS) $(CUDA_OBJS) c3_fmt.o x86.o x86-sse.o" \
 		CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \
-		LDFLAGS="$(LDFLAGS) -lcrypt"
+		LDFLAGS="$(LDFLAGS) -L$(CUDAPATH) -lcudart -lcrypt"
 
 linux-x86-mmx:
 	$(LN) x86-mmx.h arch.h
 	$(MAKE) $(PROJ) \
-		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86.o x86-mmx.o" \
+		JOHN_OBJS="$(JOHN_OBJS) $(CUDA_OBJS) c3_fmt.o x86.o x86-mmx.o" \
 		CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \
-		LDFLAGS="$(LDFLAGS) -lcrypt"
+		LDFLAGS="$(LDFLAGS) -L$(CUDAPATH) -lcudart -lcrypt"
 
 linux-x86-any:
 	$(LN) x86-any.h arch.h
 	$(MAKE) $(PROJ) \
-		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86.o" \
+		JOHN_OBJS="$(JOHN_OBJS) $(CUDA_OBJS) c3_fmt.o x86.o" \
 		CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \
-		LDFLAGS="$(LDFLAGS) -lcrypt"
+		LDFLAGS="$(LDFLAGS) -L$(CUDAPATH) -lcudart -lcrypt"
 
 linux-x86-avx:
 	$(LN) x86-sse.h arch.h
@@ -811,6 +818,13 @@ generic.h:
 		"$(BENCH_MD5_OBJS_DEPEND)" \
 		"$(BENCH_BF_OBJS_DEPEND)"
 
+cryptmd5.o:  cryptmd5.h cuda/cryptmd5.cu
+	$(shell  cd cuda; nvcc $(NVCCFLAGS) cryptmd5.cu)
+	cp cuda/cryptmd5.o cryptmd5.o
+
+cryptmd5_fmt.o: cryptmd5.o cryptmd5cuda_fmt.c cryptmd5.h
+	$(CC)  $(CFLAGS) cryptmd5cuda_fmt.c
+
 bench: $(BENCH_OBJS)
 	$(LD) $(BENCH_OBJS) $(LDFLAGS) -o bench
 
@@ -890,6 +904,8 @@ depend:
 
 clean:
 	$(RM) $(PROJ) $(PROJ_DOS) $(PROJ_WIN32)
+	$(RM) cuda/*.o cuda/*.c~ 
+	$(RM) *.cubin *.c~ *~
 	$(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/cryptmd5.h john-1.7.8-cryptmd5_7//src/cryptmd5.h
--- john-1.7.8.orig//src/cryptmd5.h	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.8-cryptmd5_7//src/cryptmd5.h	2011-07-25 05:39:19.030813316 +0000
@@ -0,0 +1,108 @@
+/*
+* This software is Copyright (c) 2011 Lukas Odzioba
+* <lukas dot odzioba at gmail dot com> 
+* 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
+#include <assert.h>
+#include "common.h"
+#include <stdint.h>
+#include <stdbool.h>
+
+#define uint32_t unsigned int
+#define uint8_t unsigned char
+
+#define MIN(a,b) ((a)<(b)?(a):(b))
+#define MAX(a,b) ((a)>(b)?(a):(b))
+
+#define BLOCKS 28
+#define THREADS 384
+#define KEYS_PER_CRYPT BLOCKS*THREADS
+#define PLAINTEXT_LENGTH	15
+
+typedef struct {
+	unsigned char saltlen;
+	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];
+} crypt_md5_password;
+
+typedef struct {
+	uint32_t v[4];		//128 bits
+} crypt_md5_hash;
+
+typedef struct {
+#define ctx_buffsize 64
+	uint8_t buffer[64];
+	uint32_t ALIGN;
+} md5_ctx;
+
+static const char md5_salt_prefix[] = "$1$";
+static const char apr1_salt_prefix[] = "$apr1$";
+
+#define ROTATE_LEFT(x, s) ((x << s) | (x >> (32 - s)))
+
+#define F(x, y, z) (x&y | ~x&z)
+#define G(x, y, z) (x&z | y&~z)
+
+#define H(x, y, z) (x^y^z)
+#define I(x, y, z) (y^(x|~z))
+
+#define FF(v, w, x, y, z, s, ac) { \
+ v += F(w, x, y) + z + ac; \
+ v = ROTATE_LEFT(v, s) + w; \
+ }
+#define GG(v, w, x, y, z, s, ac) { \
+ v += G(w, x, y) + z + ac; \
+ v = ROTATE_LEFT(v, s) + w; \
+ }
+#define HH(v, w, x, y, z, s, ac) { \
+ v += H(w, x, y) + z + ac; \
+ v = ROTATE_LEFT(v, s) + w; \
+ }
+#define II(v, w, x, y, z, s, ac) { \
+ v += I(w, x, y) + z + ac; \
+ v = ROTATE_LEFT(v, s) + w; \
+ }
+ 
+#define FF2(v, w, x, y, s, ac) { \
+ v += F(w, x, y) + ac; \
+ v = ROTATE_LEFT(v, s) + w; \
+ }
+#define GG2(v, w, x, y, s, ac) { \
+ v += G(w, x, y) +  ac; \
+ v = ROTATE_LEFT(v, s) + w; \
+ }
+#define HH2(v, w, x, y, s, ac) { \
+ v += H(w, x, y) + ac; \
+ v = ROTATE_LEFT(v, s) + w; \
+ }
+#define II2(v, w, x, y, s, ac) { \
+ v += I(w, x, y) + ac; \
+ v = ROTATE_LEFT(v, s) + w; \
+ } 
+
+#define S11 7
+#define S12 12
+#define S13 17
+#define S14 22
+#define S21 5
+#define S22 9
+#define S23 14
+#define S24 20
+#define S31 4
+#define S32 11
+#define S33 16
+#define S34 23
+#define S41 6
+#define S42 10
+#define S43 15
+#define S44 21
+
+#endif
diff -urpN john-1.7.8.orig//src/cryptmd5cuda_fmt.c john-1.7.8-cryptmd5_7//src/cryptmd5cuda_fmt.c
--- john-1.7.8.orig//src/cryptmd5cuda_fmt.c	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.8-cryptmd5_7//src/cryptmd5cuda_fmt.c	2011-07-25 05:40:03.637938824 +0000
@@ -0,0 +1,325 @@
+/*
+* This software is Copyright (c) 2011 Lukas Odzioba
+* <lukas dot odzioba at gmail dot com> 
+* 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 "cryptmd5.h"
+
+#define FORMAT_LABEL		"cryptmd5cuda"
+#define FORMAT_NAME		"CRYPTMD5CUDA"
+
+#define CRYPT_TYPE		"MD5-based CRYPT"
+
+#define BENCHMARK_COMMENT	" saltlen=8,passlen=8"
+#define BENCHMARK_LENGTH	-1
+
+#define BINARY_SIZE		16
+#define SALT_SIZE		(8+1)
+#define MIN_KEYS_PER_CRYPT	KEYS_PER_CRYPT
+#define MAX_KEYS_PER_CRYPT	KEYS_PER_CRYPT
+
+void md5_crypt_gpu(crypt_md5_password *, crypt_md5_hash *, crypt_md5_salt *);
+
+static crypt_md5_password inbuffer[MAX_KEYS_PER_CRYPT];			/** plaintext ciphertexts **/
+static crypt_md5_hash outbuffer[MAX_KEYS_PER_CRYPT];			/** calculated hashes **/
+static crypt_md5_salt host_salt;					/** salt **/
+
+//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"},
+*/
+	{NULL}
+};
+
+static void init(void)
+{
+}
+
+static int valid(char *ciphertext)
+{
+	uint8_t i, len = strlen(ciphertext), prefix = 0;
+
+	if (strncmp(ciphertext, md5_salt_prefix, strlen(md5_salt_prefix)) == 0)
+		prefix |= 1;
+	if (strncmp(ciphertext, apr1_salt_prefix,
+		strlen(apr1_salt_prefix)) == 0)
+		prefix |= 2;
+	if (prefix == 0)
+		return 0;
+
+	char *p = strrchr(ciphertext, '$');
+	for (i = p - ciphertext + 1; i < len; i++) {
+		uint8_t z = ARCH_INDEX(ciphertext[i]);
+		if (ARCH_INDEX(atoi64[z]) == 0x7f)
+			return 0;
+	}
+	if (len - (p - ciphertext + 1) != 22)
+		return 0;
+	return 1;
+};
+
+static int findb64(char c)
+{
+	int ret = ARCH_INDEX(atoi64[(uint8_t) c]);
+	return ret != 0x7f ? ret : 0;
+}
+
+static void to_binary(char *crypt, char *alt)
+{
+
+#define _24bit_from_b64(I,B2,B1,B0) \
+  do{\
+      unsigned char c1=findb64(crypt[I+0]);\
+      unsigned char c2=findb64(crypt[I+1]);\
+      unsigned char c3=findb64(crypt[I+2]);\
+      unsigned char c4=findb64(crypt[I+3]);\
+      unsigned int w=c4<<18|c3<<12|c2<<6|c1;\
+      unsigned char b2=w&0xff;w>>=8;\
+      unsigned char b1=w&0xff;w>>=8;\
+      unsigned char b0=w&0xff;w>>=8;\
+      alt[B2]=b0;\
+      alt[B1]=b1;\
+      alt[B0]=b2;\
+  }while(0)
+
+	_24bit_from_b64(0, 0, 6, 12);
+	_24bit_from_b64(4, 1, 7, 13);
+	_24bit_from_b64(8, 2, 8, 14);
+	_24bit_from_b64(12, 3, 9, 15);
+	_24bit_from_b64(16, 4, 10, 5);
+	uint32_t w = findb64(crypt[21]) << 6 | findb64(crypt[20]) << 0;
+	alt[11] = (w & 0xff);
+}
+
+static void *binary(char *ciphertext)
+{
+	static char b[BINARY_SIZE];
+	memset(b, 0, BINARY_SIZE);
+	char *p = strrchr(ciphertext, '$') + 1;
+	to_binary(p, b);
+	return (void *) b;
+}
+
+
+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;
+
+	if (strncmp(ciphertext, md5_salt_prefix, strlen(md5_salt_prefix)) == 0) {
+		pos += strlen(md5_salt_prefix);
+		ret[8] = '1';
+	}
+	if (strncmp(ciphertext, apr1_salt_prefix,
+		strlen(apr1_salt_prefix)) == 0) {
+		pos += strlen(apr1_salt_prefix);
+		ret[8] = '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 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];
+}
+
+static void set_key(char *key, int index)
+{
+	uint32_t len = strlen(key);
+	inbuffer[index].length = len;
+	memcpy((char *) inbuffer[index].v, key, len);
+}
+
+static char *get_key(int index)
+{
+	return (char *) inbuffer[index].v;
+}
+
+static void crypt_all(int count)
+{
+	md5_crypt_gpu(inbuffer, outbuffer, &host_salt);
+}
+
+static int get_hash_0(int index)
+{
+	return ((ARCH_WORD_32 *) outbuffer[index].v)[0] & 0xf;
+}
+
+static int get_hash_1(int index)
+{
+	return ((ARCH_WORD_32 *) outbuffer[index].v)[0] & 0xff;
+}
+
+static int get_hash_2(int index)
+{
+	return ((ARCH_WORD_32 *) outbuffer[index].v)[0] & 0xfff;
+}
+
+static int get_hash_3(int index)
+{
+	return ((ARCH_WORD_32 *) outbuffer[index].v)[0] & 0xffff;
+}
+
+static int get_hash_4(int index)
+{
+	return ((ARCH_WORD_32 *) 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, *t = (uint32_t *) binary;
+	for (i = 0; i < 4; i++)
+		if (t[i] != outbuffer[index].v[i])
+			return 0;
+	return 1;
+}
+
+static int cmp_exact(char *source, int count)
+{
+	return 1;
+}
+
+struct fmt_main fmt_CRYPTMD5CUDA = {
+	{
+		    FORMAT_LABEL,
+		    FORMAT_NAME,
+		    CRYPT_TYPE,
+		    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/cuda/cryptmd5.cu john-1.7.8-cryptmd5_7//src/cuda/cryptmd5.cu
--- john-1.7.8.orig//src/cuda/cryptmd5.cu	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.8-cryptmd5_7//src/cuda/cryptmd5.cu	2011-07-25 05:39:39.069057693 +0000
@@ -0,0 +1,268 @@
+/*
+* This software is Copyright (c) 2011 Lukas Odzioba
+* <lukas dot odzioba at gmail dot com> 
+* 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 <stdio.h>
+#include <stdlib.h>
+#include <assert.h>
+#include <string.h>
+#include "../cryptmd5.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 md5_crypt_gpu(crypt_md5_password *, crypt_md5_hash *,
+    crypt_md5_salt *);
+
+__device__ __constant__ char md5_salt_prefix_cu[] = "$1$";
+__device__ __constant__ char apr1_salt_prefix_cu[] = "$apr1$";
+__device__ __constant__ crypt_md5_salt cuda_salt[1];
+
+__device__ void md5_process_block_cu(const void *, size_t, md5_ctx *);
+__device__ void md5_process_bytes_cu(const void *, size_t, md5_ctx *);
+
+__device__ void init_ctx(md5_ctx * ctx,uint8_t *ctx_buflen)
+{
+	int i = (ctx_buffsize / sizeof(uint32_t))-2;
+	uint32_t *buf = (uint32_t *) ctx->buffer;
+	while (i--)
+		*buf++ = 0;
+	ctx_buflen[threadIdx.x]=0;
+}
+
+__device__ void ctx_update(md5_ctx * ctx, const char *string, uint8_t len,uint8_t *ctx_buflen)
+{
+	uint8_t *dest = &ctx->buffer[ctx_buflen[threadIdx.x]];
+	uint8_t *src = (uint8_t *) string;
+	ctx_buflen[threadIdx.x]+=len;
+	memcpy(dest,src,len);
+}
+
+__device__ void md5_digest(md5_ctx * ctx, uint32_t * result,uint8_t *ctx_buflen)
+{
+  	uint32_t len = ctx_buflen[threadIdx.x];
+	uint32_t *x = (uint32_t *) ctx->buffer;
+	x[len / 4] |= (((uint32_t) 0x80) << ((len & 0x3) << 3));
+	len<<=3;
+	  
+	uint32_t a = 0x67452301;
+	uint32_t b = 0xefcdab89;
+	uint32_t c = 0x98badcfe;
+	uint32_t d = 0x10325476;
+
+	{
+		FF(a, b, c, d, x[0], S11, 0xd76aa478);	/* 1 */
+		FF(d, a, b, c, x[1], S12, 0xe8c7b756);	/* 2 */
+		FF(c, d, a, b, x[2], S13, 0x242070db);	/* 3 */
+		FF(b, c, d, a, x[3], S14, 0xc1bdceee);	/* 4 */
+		FF(a, b, c, d, x[4], S11, 0xf57c0faf);	/* 5 */
+		FF(d, a, b, c, x[5], S12, 0x4787c62a);	/* 6 */
+		FF(c, d, a, b, x[6], S13, 0xa8304613);	/* 7 */
+		FF(b, c, d, a, x[7], S14, 0xfd469501);	/* 8 */
+		FF(a, b, c, d, x[8], S11, 0x698098d8);	/* 9 */
+		FF(d, a, b, c, x[9], S12, 0x8b44f7af);	/* 10 */
+		FF(c, d, a, b, x[10], S13, 0xffff5bb1);	/* 11 */
+		FF(b, c, d, a, x[11], S14, 0x895cd7be);	/* 12 */
+		FF(a, b, c, d, x[12], S11, 0x6b901122);	/* 13 */
+		FF(d, a, b, c, x[13], S12, 0xfd987193);	/* 14 */
+		FF(c, d, a, b, len, S13, 0xa679438e);	/* 15 */
+		FF2(b, c, d, a, S14, 0x49b40821);	/* 16 */
+	}
+	{
+		GG(a, b, c, d, x[1], S21, 0xf61e2562);	/* 17 */
+		GG(d, a, b, c, x[6], S22, 0xc040b340);	/* 18 */
+		GG(c, d, a, b, x[11], S23, 0x265e5a51);	/* 19 */
+		GG(b, c, d, a, x[0], S24, 0xe9b6c7aa);	/* 20 */
+		GG(a, b, c, d, x[5], S21, 0xd62f105d);	/* 21 */
+		GG(d, a, b, c, x[10], S22, 0x2441453);	/* 22 */
+		GG2(c, d, a, b, S23, 0xd8a1e681);	/* 23 */
+		GG(b, c, d, a, x[4], S24, 0xe7d3fbc8);	/* 24 */
+		GG(a, b, c, d, x[9], S21, 0x21e1cde6);	/* 25 */
+		GG(d, a, b, c,len, S22, 0xc33707d6);	/* 26 */
+		GG(c, d, a, b, x[3], S23, 0xf4d50d87);	/* 27 */
+		GG(b, c, d, a, x[8], S24, 0x455a14ed);	/* 28 */
+		GG(a, b, c, d, x[13], S21, 0xa9e3e905);	/* 29 */
+		GG(d, a, b, c, x[2], S22, 0xfcefa3f8);	/* 30 */
+		GG(c, d, a, b, x[7], S23, 0x676f02d9);	/* 31 */
+		GG(b, c, d, a, x[12], S24, 0x8d2a4c8a);	/* 32 */
+	}
+	{
+		HH(a, b, c, d, x[5], S31, 0xfffa3942);	/* 33 */
+		HH(d, a, b, c, x[8], S32, 0x8771f681);	/* 34 */
+		HH(c, d, a, b, x[11], S33, 0x6d9d6122);	/* 35 */
+		HH(b, c, d, a, len, S34, 0xfde5380c);	/* 36 */
+		HH(a, b, c, d, x[1], S31, 0xa4beea44);	/* 37 */
+		HH(d, a, b, c, x[4], S32, 0x4bdecfa9);	/* 38 */
+		HH(c, d, a, b, x[7], S33, 0xf6bb4b60);	/* 39 */
+		HH(b, c, d, a, x[10], S34, 0xbebfbc70);	/* 40 */
+		HH(a, b, c, d, x[13], S31, 0x289b7ec6);	/* 41 */
+		HH(d, a, b, c, x[0], S32, 0xeaa127fa);	/* 42 */
+		HH(c, d, a, b, x[3], S33, 0xd4ef3085);	/* 43 */
+		HH(b, c, d, a, x[6], S34, 0x4881d05);	/* 44 */
+		HH(a, b, c, d, x[9], S31, 0xd9d4d039);	/* 45 */
+		HH(d, a, b, c, x[12], S32, 0xe6db99e5);	/* 46 */
+		HH2(c, d, a, b, S33, 0x1fa27cf8);	/* 47 */
+		HH(b, c, d, a, x[2], S34, 0xc4ac5665);	/* 48 */
+	}
+	{
+		II(a, b, c, d, x[0], S41, 0xf4292244);	/* 49 */
+		II(d, a, b, c, x[7], S42, 0x432aff97);	/* 50 */
+		II(c, d, a, b, len, S43, 0xab9423a7);	/* 51 */
+		II(b, c, d, a, x[5], S44, 0xfc93a039);	/* 52 */
+		II(a, b, c, d, x[12], S41, 0x655b59c3);	/* 53 */
+		II(d, a, b, c, x[3], S42, 0x8f0ccc92);	/* 54 */
+		II(c, d, a, b, x[10], S43, 0xffeff47d);	/* 55 */
+		II(b, c, d, a, x[1], S44, 0x85845dd1);	/* 56 */
+		II(a, b, c, d, x[8], S41, 0x6fa87e4f);	/* 57 */
+		II2(d, a, b, c, S42, 0xfe2ce6e0);	/* 58 */
+		II(c, d, a, b, x[6], S43, 0xa3014314);	/* 59 */
+		II(b, c, d, a, x[13], S44, 0x4e0811a1);	/* 60 */
+		II(a, b, c, d, x[4], S41, 0xf7537e82);	/* 61 */
+		II(d, a, b, c, x[11], S42, 0xbd3af235);	/* 62 */
+		II(c, d, a, b, x[2], S43, 0x2ad7d2bb);	/* 63 */
+		II(b, c, d, a, x[9], S44, 0xeb86d391);	/* 64 */
+	}
+	a += 0x67452301;
+	b += 0xefcdab89;
+	c += 0x98badcfe;
+	d += 0x10325476;
+
+	result[0]=a;
+	result[1]=b;
+	result[2]=c;
+	result[3]=d;
+}
+
+
+__device__ void md5crypt(const char *gpass, size_t keysize, uint32_t * tresult,
+    uint32_t idx)
+{
+
+	uint32_t i;
+	__shared__ uint32_t alt_result[THREADS][4 + 1];
+	__shared__ char spass[THREADS][16 + 4];
+	__shared__ uint8_t ctx_buflen[THREADS];
+	__shared__ uint8_t altctx_buflen[THREADS];
+	
+
+	char *pass = spass[threadIdx.x];
+	memcpy(pass, gpass, 16);
+	uint8_t pass_len = keysize;
+	uint8_t salt_len = cuda_salt[0].saltlen;
+	char *salt = cuda_salt[0].salt;
+	md5_ctx ctx, alt_ctx;
+	init_ctx(&ctx,ctx_buflen);
+	init_ctx(&alt_ctx,altctx_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);
+	} else
+		ctx_update(&ctx, apr1_salt_prefix_cu, 6,ctx_buflen);
+
+
+	ctx_update(&ctx, salt, salt_len,ctx_buflen);
+
+
+	ctx_update(&alt_ctx, pass, pass_len,altctx_buflen);
+	ctx_update(&alt_ctx, salt, salt_len,altctx_buflen);
+	ctx_update(&alt_ctx, pass, pass_len,altctx_buflen);
+	md5_digest(&alt_ctx, alt_result[threadIdx.x],altctx_buflen);
+
+	for (i = pass_len; i > 16; i -= 16)
+		ctx_update(&ctx, (const char *) alt_result[threadIdx.x], 16,ctx_buflen);
+	ctx_update(&ctx, (const char *) alt_result[threadIdx.x], i,ctx_buflen);
+
+	*alt_result[threadIdx.x] = 0;
+
+	for (i = pass_len; i > 0; i >>= 1)
+		if ((i & 1) != 0)
+			ctx.buffer[ctx_buflen[threadIdx.x]++] =
+			    ((const char *) alt_result[threadIdx.x])[0];
+		else
+			ctx.buffer[ctx_buflen[threadIdx.x]++] = pass[0];
+
+	md5_digest(&ctx, alt_result[threadIdx.x],ctx_buflen);
+
+	for (i = 0; i < 1000; i++) {
+		init_ctx(&ctx,ctx_buflen);
+
+		if ((i & 1) != 0)
+			ctx_update(&ctx, pass, pass_len,ctx_buflen);
+		else
+			ctx_update(&ctx,
+			    (const char *) alt_result[threadIdx.x], 16,ctx_buflen);
+
+		if (i % 3 != 0)
+			ctx_update(&ctx, salt, salt_len,ctx_buflen);
+
+		if (i % 7 != 0)
+			ctx_update(&ctx, pass, pass_len,ctx_buflen);
+
+		if ((i & 1) != 0)
+			ctx_update(&ctx,
+			    (const char *) alt_result[threadIdx.x], 16,ctx_buflen);
+		else
+			ctx_update(&ctx, pass, pass_len,ctx_buflen);
+		md5_digest(&ctx, alt_result[threadIdx.x],ctx_buflen);
+
+	}
+	//__syncthreads();
+	tresult[0] = alt_result[threadIdx.x][0];
+	tresult[1] = alt_result[threadIdx.x][1];
+	tresult[2] = alt_result[threadIdx.x][2];
+	tresult[3] = alt_result[threadIdx.x][3];
+	//__syncthreads();
+
+}
+
+
+__global__ void kernel_crypt_r(crypt_md5_password * inbuffer,
+    crypt_md5_hash * outbuffer)
+{
+	uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+	md5crypt((char *) inbuffer[idx].v, inbuffer[idx].length,
+	    outbuffer[idx].v, idx);
+}
+
+__host__ void md5_crypt_gpu(crypt_md5_password * inbuffer,
+    crypt_md5_hash * outbuffer, crypt_md5_salt * host_salt)
+{
+
+
+	HANDLE_ERROR(cudaMemcpyToSymbol(cuda_salt, host_salt,
+		sizeof(crypt_md5_salt)));
+	crypt_md5_password *cuda_inbuffer;
+	crypt_md5_hash *cuda_outbuffer;
+
+	size_t insize = sizeof(crypt_md5_password) * KEYS_PER_CRYPT;
+	size_t outsize = sizeof(crypt_md5_hash) * KEYS_PER_CRYPT;
+
+	HANDLE_ERROR(cudaMalloc(&cuda_inbuffer, insize));
+	HANDLE_ERROR(cudaMalloc(&cuda_outbuffer, outsize));
+	HANDLE_ERROR(cudaMemcpy(cuda_inbuffer, inbuffer, insize,
+		cudaMemcpyHostToDevice));
+	dim3 dimGrid(BLOCKS);
+	dim3 dimBlock(THREADS);
+
+	kernel_crypt_r <<< dimGrid, dimBlock >>> (cuda_inbuffer,
+	    cuda_outbuffer);
+
+	cudaThreadSynchronize();
+
+	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-cryptmd5_7//src/john.c
--- john-1.7.8.orig//src/john.c	2011-02-27 12:31:36.000000000 +0000
+++ john-1.7.8-cryptmd5_7//src/john.c	2011-07-25 05:32:23.374813746 +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_CRYPTMD5CUDA;
 #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_CRYPTMD5CUDA);
 #ifdef HAVE_CRYPT
 	john_register_one(&fmt_crypt);
 #endif
diff -urpN john-1.7.8.orig//src/options.c john-1.7.8-cryptmd5_7//src/options.c
--- john-1.7.8.orig//src/options.c	2011-06-22 13:01:40.000000000 +0000
+++ john-1.7.8-cryptmd5_7//src/options.c	2011-07-25 05:31:34.207063739 +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/CRYPTMD5CUDA" MAYBE_CRYPT "\n" \
 "--save-memory=LEVEL        enable memory saving, at LEVEL 1..3\n"
 
 void opt_init(char *name, int argc, char **argv)
