diff -urpN john-1.7.8.orig//doc/PHPASS_CUDA john-1.7.8-phpasscuda_5//doc/PHPASS_CUDA
--- john-1.7.8.orig//doc/PHPASS_CUDA	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.8-phpasscuda_5//doc/PHPASS_CUDA	2011-08-03 23:00:16.808488625 +0000
@@ -0,0 +1,80 @@
+  ,-*
+ (_)   Release Notes
+
+==== Version        ====
+    PHPASS CUDA ver. 3.0 patch for John The Ripper 1.7.8
+
+==== Release Date   ====
+    DD.MM.RRRR = 04.08.2011
+
+==== Authors        ====
+    Lukas Odzioba, lukas[dot]odzioba[at]gmail[dot]com
+
+==== Change Log     ====
+      ver 3:
+	+22  speedup
+      ver 2:
+	+25% speedup
+      ver 1:
+	LoopCnt is now kernel argument.
+	Standarized tests to $P$9* prefix
+	Removed validchar() function
+	Major speed improvements, increased performance by 85%
+	Improved this document, moved to doc, renamed
+      ver 0:
+	First revision
+       
+==== Restrictions   ====
+    Patch was tested/optimized for following configuration:
+    i3-2100, 4GB 1333Mhz, Slackware 13.1, GTX 460
+    Patch is not setting up THREADS and BLOCKS automaticaly, you should experiment with those constants 
+    to get better performance on your card.
+    Passwords length less than 16.
+
+==== ToDo 	    ====
+    Add $H$ support
+    Add an autoconfiguration script
+    Further speed improvements
+
+==== Known Bugs     ====
+    Patch is having trouble with long execution of grater phpass loop counters.
+    Executions is, aborted by OS. Changing BLOCKS to smaller should partially help.
+
+==== Environment    ====
+    PHPASSCUDA is supported/tested only on following builds:
+      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    ====
+    Apply this patch on top of JtR 1.7.8
+    $ cd john-1.7.8
+    $ patch -p1 < ../john-1.7.8-phpasscuda-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          ====
+    Incremental attack:
+    $ john-1.7.8/run/john -i -format=phpasscuda FILE
+
+==== Benchmarks     ==== 
+    ver 0:
+      GeForce 9800GT			Raw:  27459 c/s real,  27510 c/s virtual
+      GeForce 460m			Raw:  49461 c/s real,  49461 c/s virtual
+      GeForce GTX 460 1024MB		Raw: 102400 c/s real, 102603 c/s virtual
+      GeForce GTX 570			Raw: 177075 c/s real, 177075 c/s virtual
+    ver 1:
+      GeForce GTX 460 1024MB		Raw: 187200 c/s real, 187560 c/s virtual
+    ver 2:
+      GeForce GTX 460 1024MB		Raw: 237804 c/s real, 237804 c/s virtual
+    ver 3:
+      GeForce GTX 460 1024MB		Raw: 292129 c/s real, 292129 c/s virtual
+
+
diff -urpN john-1.7.8.orig//src/Makefile john-1.7.8-phpasscuda_5//src/Makefile
--- john-1.7.8.orig//src/Makefile	2011-05-04 18:52:31.000000000 +0000
+++ john-1.7.8-phpasscuda_5//src/Makefile	2011-07-25 08:55:50.000000000 +0000
@@ -3,6 +3,7 @@
 # Copyright (c) 1996-2011 by Solar Designer
 #
 
+NVCC = nvcc
 CC = gcc
 AS = $(CC)
 LD = $(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_10 --maxrregcount=24
+CUDAPATH = /usr/local/cuda/lib
+CUDA64PATH = /usr/local/cuda/lib64
 CFLAGS = -c -Wall -O2 -fomit-frame-pointer $(OMPFLAGS)
 ASFLAGS = -c $(OMPFLAGS)
 LDFLAGS = -s $(OMPFLAGS)
@@ -33,6 +37,7 @@ JOHN_OBJS = \
 	AFS_fmt.o \
 	LM_fmt.o \
 	dummy.o \
+	phpasscuda_fmt.o phpass.o \
 	batch.o bench.o charset.o common.o compiler.o config.o cracker.o \
 	crc32.o external.o formats.o getopt.o idle.o inc.o john.o list.o \
 	loader.o logger.o math.o memory.o misc.o options.o params.o path.o \
@@ -153,7 +158,7 @@ linux-x86-64:
 	$(MAKE) $(PROJ) \
 		JOHN_OBJS="$(JOHN_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
@@ -190,21 +195,21 @@ linux-x86-sse2:
 	$(MAKE) $(PROJ) \
 		JOHN_OBJS="$(JOHN_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" \
 		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" \
 		CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \
-		LDFLAGS="$(LDFLAGS) -lcrypt"
+		LDFLAGS="$(LDFLAGS) -L$(CUDAPATH) -lcudart -lcrypt"
 
 linux-x86-avx:
 	$(LN) x86-sse.h arch.h
@@ -811,6 +816,13 @@ generic.h:
 		"$(BENCH_MD5_OBJS_DEPEND)" \
 		"$(BENCH_BF_OBJS_DEPEND)"
 
+phpass.o:  phpass.h cuda/phpass.cu
+	$(shell  cd cuda; nvcc $(NVCCFLAGS) phpass.cu)
+
+phpasscuda_fmt.o: phpass.o phpasscuda_fmt.c phpass.h
+	cp cuda/phpass.o phpass.o
+	$(CC)  $(CFLAGS) phpasscuda_fmt.c
+
 bench: $(BENCH_OBJS)
 	$(LD) $(BENCH_OBJS) $(LDFLAGS) -o bench
 
@@ -890,6 +902,7 @@ depend:
 
 clean:
 	$(RM) $(PROJ) $(PROJ_DOS) $(PROJ_WIN32)
+	$(RM) cuda/*.o cuda/*~ *~
 	$(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/phpass.cu john-1.7.8-phpasscuda_5//src/cuda/phpass.cu
--- john-1.7.8.orig//src/cuda/phpass.cu	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.8-phpasscuda_5//src/cuda/phpass.cu	2011-08-03 22:53:49.406575801 +0000
@@ -0,0 +1,224 @@
+/*
+* 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 "../phpass.h"
+
+const uint SETTING_SIZE = 12;
+const uint DATA_IN_SIZE = KEYS_PER_CRYPT * sizeof(phpass_password);
+const uint DATA_OUT_SIZE = KEYS_PER_CRYPT * sizeof(phpass_hash);
+
+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__))
+
+__device__ unsigned char *cuda_data = NULL;
+__device__ uint32_t *cuda_data_out = NULL;
+__device__ char *cuda_setting = NULL;
+
+unsigned char *host_data = NULL;
+uint32_t *host_data_out = NULL;
+
+extern "C" void mem_init(unsigned char *, uint32_t *, char *, char *, int);
+extern "C" void mem_clear(void);
+extern "C" void gpu_phpass(void);
+
+__global__ void kernel_phpass(unsigned char *data, uint32_t * data_out,
+    char *, int);
+
+
+int cuda_count_log2;
+__host__ void gpu_phpass(void)
+{
+	dim3 dimGrid(BLOCKS);
+	dim3 dimBlock(THREADS);
+	kernel_phpass <<< dimGrid, dimBlock >>> (cuda_data, cuda_data_out,
+	    cuda_setting, cuda_count_log2);
+	HANDLE_ERROR(cudaThreadSynchronize());
+}
+
+__host__ void mem_init(unsigned char *p, uint32_t * h, char *setting,
+    char *itoa, int count_log2)
+{
+	cuda_count_log2 = count_log2;
+	host_data = p;
+	host_data_out = h;
+	HANDLE_ERROR(cudaMalloc(&cuda_setting, SETTING_SIZE));
+	HANDLE_ERROR(cudaMalloc(&cuda_data, DATA_IN_SIZE));
+	HANDLE_ERROR(cudaMalloc(&cuda_data_out, DATA_OUT_SIZE));
+	HANDLE_ERROR(cudaMemcpy(cuda_data, host_data, DATA_IN_SIZE,
+		cudaMemcpyHostToDevice));
+	HANDLE_ERROR(cudaMemcpy(cuda_setting, setting, SETTING_SIZE,
+		cudaMemcpyHostToDevice));
+}
+
+__host__ void mem_clear()
+{
+	HANDLE_ERROR(cudaMemcpy(host_data_out, cuda_data_out, DATA_OUT_SIZE,
+		cudaMemcpyDeviceToHost));
+	HANDLE_ERROR(cudaFree(cuda_data));
+	HANDLE_ERROR(cudaFree(cuda_data_out));
+	HANDLE_ERROR(cudaFree(cuda_setting));
+}
+
+__device__ void cuda_md5(char len, uint32_t * internal_ret, uint32_t * x)
+{
+	x[len / 4] |= (((uint32_t) 0x80) << ((len & 0x3) << 3));
+	uint32_t x14 = len << 3;
+
+	uint32_t a = 0x67452301;
+	uint32_t b = 0xefcdab89;
+	uint32_t c = 0x98badcfe;
+	uint32_t d = 0x10325476;
+
+	a = AC1 + x[0];
+	a = ROTATE_LEFT(a, S11);
+	a += b;			/* 1 */
+	d = (c ^ (a & MASK1)) + x[1] + AC2pCd;
+	d = ROTATE_LEFT(d, S12);
+	d += a;			/* 2 */
+	c = F(d, a, b) + x[2] + AC3pCc;
+	c = ROTATE_LEFT(c, S13);
+	c += d;			/* 3 */
+	b = F(c, d, a) + x[3] + AC4pCb;
+	b = ROTATE_LEFT(b, S14);
+	b += c;			/* 4 */
+	FF(a, b, c, d, x[4], S11, 0xf57c0faf);
+	FF(d, a, b, c, x[5], S12, 0x4787c62a);
+	FF(c, d, a, b, x[6], S13, 0xa8304613);
+	FF(b, c, d, a, x[7], S14, 0xfd469501);
+	FF(a, b, c, d, x[8], S11, 0x698098d8);
+	FF(d, a, b, c, x[9], S12, 0x8b44f7af);
+	FF(c, d, a, b, x[10], S13, 0xffff5bb1);
+	FF(b, c, d, a, x[11], S14, 0x895cd7be);
+	FF(a, b, c, d, x[12], S11, 0x6b901122);
+	FF(d, a, b, c, x[13], S12, 0xfd987193);
+	FF(c, d, a, b, x14, S13, 0xa679438e);
+	FF(b, c, d, a, 0, S14, 0x49b40821);
+
+	GG(a, b, c, d, x[1], S21, 0xf61e2562);
+	GG(d, a, b, c, x[6], S22, 0xc040b340);
+	GG(c, d, a, b, x[11], S23, 0x265e5a51);
+	GG(b, c, d, a, x[0], S24, 0xe9b6c7aa);
+	GG(a, b, c, d, x[5], S21, 0xd62f105d);
+	GG(d, a, b, c, x[10], S22, 0x2441453);
+	GG(c, d, a, b, 0, S23, 0xd8a1e681);
+	GG(b, c, d, a, x[4], S24, 0xe7d3fbc8);
+	GG(a, b, c, d, x[9], S21, 0x21e1cde6);
+	GG(d, a, b, c, x14, S22, 0xc33707d6);
+	GG(c, d, a, b, x[3], S23, 0xf4d50d87);
+	GG(b, c, d, a, x[8], S24, 0x455a14ed);
+	GG(a, b, c, d, x[13], S21, 0xa9e3e905);
+	GG(d, a, b, c, x[2], S22, 0xfcefa3f8);
+	GG(c, d, a, b, x[7], S23, 0x676f02d9);
+	GG(b, c, d, a, x[12], S24, 0x8d2a4c8a);
+
+	HH(a, b, c, d, x[5], S31, 0xfffa3942);
+	HH(d, a, b, c, x[8], S32, 0x8771f681);
+	HH(c, d, a, b, x[11], S33, 0x6d9d6122);
+	HH(b, c, d, a, x14, S34, 0xfde5380c);
+	HH(a, b, c, d, x[1], S31, 0xa4beea44);
+	HH(d, a, b, c, x[4], S32, 0x4bdecfa9);
+	HH(c, d, a, b, x[7], S33, 0xf6bb4b60);
+	HH(b, c, d, a, x[10], S34, 0xbebfbc70);
+	HH(a, b, c, d, x[13], S31, 0x289b7ec6);
+	HH(d, a, b, c, x[0], S32, 0xeaa127fa);
+	HH(c, d, a, b, x[3], S33, 0xd4ef3085);
+	HH(b, c, d, a, x[6], S34, 0x4881d05);
+	HH(a, b, c, d, x[9], S31, 0xd9d4d039);
+	HH(d, a, b, c, x[12], S32, 0xe6db99e5);
+	HH(c, d, a, b, 0, S33, 0x1fa27cf8);
+	HH(b, c, d, a, x[2], S34, 0xc4ac5665);
+
+	II(a, b, c, d, x[0], S41, 0xf4292244);
+	II(d, a, b, c, x[7], S42, 0x432aff97);
+	II(c, d, a, b, x14, S43, 0xab9423a7);
+	II(b, c, d, a, x[5], S44, 0xfc93a039);
+	II(a, b, c, d, x[12], S41, 0x655b59c3);
+	II(d, a, b, c, x[3], S42, 0x8f0ccc92);
+	II(c, d, a, b, x[10], S43, 0xffeff47d);
+	II(b, c, d, a, x[1], S44, 0x85845dd1);
+	II(a, b, c, d, x[8], S41, 0x6fa87e4f);
+	II(d, a, b, c, 0, S42, 0xfe2ce6e0);
+	II(c, d, a, b, x[6], S43, 0xa3014314);
+	II(b, c, d, a, x[13], S44, 0x4e0811a1);
+	II(a, b, c, d, x[4], S41, 0xf7537e82);
+	II(d, a, b, c, x[11], S42, 0xbd3af235);
+	II(c, d, a, b, x[2], S43, 0x2ad7d2bb);
+	II(b, c, d, a, x[9], S44, 0xeb86d391);
+
+	internal_ret[0] = a + 0x67452301;
+	internal_ret[1] = b + 0xefcdab89;
+	internal_ret[2] = c + 0x98badcfe;
+	internal_ret[3] = d + 0x10325476;
+}
+
+__device__ void clear_ctx(uint32_t * x)
+{
+	int i;
+#pragma unroll 12
+	for (i = 2; i < 14; i++)
+		*x++ = 0;
+}
+
+__device__ void clear_ctx_fast(uint32_t * x)
+{
+	int i;
+#pragma unroll 10
+	for (i = 4; i < 14; i++)
+		*x++ = 0;
+}
+
+__global__ void kernel_phpass(unsigned char *password, uint32_t * data_out,
+    char *setting, int count_log2)
+{
+	uint32_t ret[4];
+	uint32_t x[14];
+	int length, count, i;
+	unsigned char *buff = (unsigned char *) x;
+	clear_ctx(x);
+	uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+
+	length = password[address(15, idx)];
+
+#pragma unroll 8
+	for (i = 0; i < 8; i++)
+		buff[i] = setting[i + 4];
+
+	for (i = 8; i < 8 + length; i++) {
+		buff[i] = password[address(i - 8, idx)];
+	}
+	
+	cuda_md5(8 + length, ret, x);
+	count = 1 << count_log2;
+	clear_ctx_fast(x);
+	for (i = 16; i < 16 + length; i++)
+		buff[i] = password[address(i - 16, idx)];
+
+#pragma unroll 4
+	for (i = 0; i < 4; i++)
+		x[i] = ret[i];
+	uint32_t *ret2 = x;
+	do {
+		cuda_md5(16 + length, ret2, x);
+
+	} while (--count);
+
+#pragma unroll 4
+	for (i = 0; i < 4; i++)
+		data_out[address(i, idx)] = x[i];
+
+}
diff -urpN john-1.7.8.orig//src/john.c john-1.7.8-phpasscuda_5//src/john.c
--- john-1.7.8.orig//src/john.c	2011-02-27 12:31:36.000000000 +0000
+++ john-1.7.8-phpasscuda_5//src/john.c	2011-07-25 08:52:53.000000000 +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_PHPASSCUDA;
 #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_PHPASSCUDA);
 #ifdef HAVE_CRYPT
 	john_register_one(&fmt_crypt);
 #endif
diff -urpN john-1.7.8.orig//src/options.c john-1.7.8-phpasscuda_5//src/options.c
--- john-1.7.8.orig//src/options.c	2011-06-22 13:01:40.000000000 +0000
+++ john-1.7.8-phpasscuda_5//src/options.c	2011-07-25 08:52:22.000000000 +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/PHPASSCUDA" MAYBE_CRYPT "\n" \
 "--save-memory=LEVEL        enable memory saving, at LEVEL 1..3\n"
 
 void opt_init(char *name, int argc, char **argv)
diff -urpN john-1.7.8.orig//src/phpass.h john-1.7.8-phpasscuda_5//src/phpass.h
--- john-1.7.8.orig//src/phpass.h	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.8-phpasscuda_5//src/phpass.h	2011-08-03 22:55:34.917518838 +0000
@@ -0,0 +1,81 @@
+/*
+* 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 _PHPASS_H
+#define _PHPASS_H
+#include "common.h"
+
+#define uint8_t 			unsigned char
+#define uint32_t 			unsigned int
+
+#define ROTATE_LEFT(x, s) 		((x << s) | (x >> (32 - s)))
+#define BLOCKS 				126*3	//it must be always something*3
+#define THREADS 			256
+#define KEYS_PER_CRYPT 			BLOCKS*THREADS
+
+#define F(x, y, z)			((z) ^ ((x) & ((y) ^ (z))))
+#define G(x, y, z)			((y) ^ ((z) & ((x) ^ (y))))
+#define H(x, y, z)			((x) ^ (y) ^ (z))
+#define I(x, y, z)			((y) ^ ((x) | ~(z)))
+
+#define address(j,idx) 			(((j)*KEYS_PER_CRYPT)+(idx))
+
+#define FF(a, b, c, d, x, s, ac) \
+  {(a) += F ((b), (c), (d)) + (x) + (uint32_t)(ac); \
+   (a) = ROTATE_LEFT ((a), (s)); \
+   (a) += (b); \
+  }
+#define GG(a, b, c, d, x, s, ac) \
+  {(a) += G ((b), (c), (d)) + (x) + (uint32_t)(ac); \
+   (a) = ROTATE_LEFT ((a), (s)); \
+   (a) += (b); \
+  }
+#define HH(a, b, c, d, x, s, ac) \
+  {(a) += H ((b), (c), (d)) + (x) + (uint32_t)(ac); \
+   (a) = ROTATE_LEFT ((a), (s)); \
+   (a) += (b); \
+  }
+#define II(a, b, c, d, x, s, ac) \
+  {(a) += I ((b), (c), (d)) + (x) + (uint32_t)(ac); \
+   (a) = ROTATE_LEFT ((a), (s)); \
+   (a) += (b); \
+  }
+
+#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
+
+#define AC1				0xd76aa477
+#define AC2pCd				0xf8fa0bcc
+#define AC3pCc				0xbcdb4dd9
+#define AC4pCb				0xb18b7a77
+#define MASK1				0x77777777
+
+static char phpass_prefix[] = "$P$";
+
+typedef struct {
+	unsigned char v[15];	///more than we need but faster
+	unsigned char length;
+} phpass_password;
+
+typedef struct {
+	uint32_t v[4];		///128bits for hash
+} phpass_hash;
+
+#endif
diff -urpN john-1.7.8.orig//src/phpasscuda_fmt.c john-1.7.8-phpasscuda_5//src/phpasscuda_fmt.c
--- john-1.7.8.orig//src/phpasscuda_fmt.c	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.8-phpasscuda_5//src/phpasscuda_fmt.c	2011-08-03 22:56:41.631738618 +0000
@@ -0,0 +1,294 @@
+/*
+* 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 "arch.h"
+#include "formats.h"
+#include "common.h"
+#include "misc.h"
+#include "phpass.h"
+
+#define FORMAT_LABEL		"phpasscuda"
+#define FORMAT_NAME		"PHPASSCUDA"
+
+#define PHPASS_TYPE		"PORTABLE-MD5"
+
+#define BENCHMARK_COMMENT	""
+#define BENCHMARK_LENGTH	-1
+
+#define PLAINTEXT_LENGTH	15
+#define CIPHERTEXT_LENGTH	34	/// header = 3 | loopcnt = 1 | salt = 8 | ciphertext = 22
+#define BINARY_SIZE		16
+#define MD5_DIGEST_LENGTH 	16
+
+
+#define SALT_SIZE		8
+
+#define MIN_KEYS_PER_CRYPT	KEYS_PER_CRYPT
+#define MAX_KEYS_PER_CRYPT	KEYS_PER_CRYPT
+
+static unsigned char inbuffer[MAX_KEYS_PER_CRYPT * sizeof(phpass_password)];			/** plaintext ciphertexts **/
+static uint32_t outbuffer[MAX_KEYS_PER_CRYPT * 4];						/** calculated hashes **/
+
+static char currentsalt[SALT_SIZE];
+static char loopChar = '*';
+
+extern void mem_init(unsigned char *, uint32_t *, char *, char *, int);
+extern void mem_clear(void);
+extern void gpu_phpass(void);
+
+static struct fmt_tests tests[] = {
+	{"$P$900000000jPBDh/JWJIyrF0.DmP7kT.", "ala"},
+	{"$P$900000000a94rg7R/nUK0icmALICKj1", "john"},
+	{"$P$900000001ahWiA6cMRZxkgUxj4x/In0", "john"},
+	{"$P$900000000m6YEJzWtTmNBBL4jypbHv1", "openwall"},
+	{"$P$900000000zgzuX4Dc2091D8kak8RdR0", "h3ll00"},
+	{"$P$900000000qZTL5A0XQUX9hq0t8SoKE0", "1234567890"},
+	{"$P$900112200B9LMtPy2FSq910c1a6BrH0", "1234567890"},
+	{"$P$900000000a94rg7R/nUK0icmALICKj1", "john"},
+	{"$P$9sadli2.wzQIuzsR2nYVhUSlHNKgG/0", "john"},
+	{"$P$90000000000tbNYOc9TwXvLEI62rPt1", ""},
+
+//      {"$P$9saltstriAcRMGl.91RgbAD6WSq64z.","a"},
+//      {"$P$9saltstriMljTzvdluiefEfDeGGQEl/","ab"},
+//      {"$P$9saltstrikCftjZCE7EY2Kg/pjbl8S.","abc"},
+//      {"$P$9saltstriV/GXRIRi9UVeMLMph9BxF0","abcd"},
+//      {"$P$9saltstri3JPgLni16rBZtI03oeqT.0","abcde"},
+//      {"$P$9saltstri0D3A6JyITCuY72ZoXdejV.","abcdef"},
+//      {"$P$9saltstriXeNc.xV8N.K9cTs/XEn13.","abcdefg"},
+//      {"$P$9saltstrinwvfzVRP3u1gxG2gTLWqv.","abcdefgh"},
+//      {"$P$9saltstriSUQTD.yC2WigjF8RU0Q.Z.","abcdefghi"},
+//      {"$P$9saltstriWPpGLG.jwJkwGRwdKNEsg.","abcdefghij"},
+
+//      {"$P$9saltstrizjDEWUMXTlQHQ3/jhpR4C.","abcdefghijk"},
+//      {"$P$9saltstriGLUwnE6bl91BPJP6sxyka.","abcdefghijkl"},
+//      {"$P$9saltstriq7s97e2m7dXnTEx2mtPzx.","abcdefghijklm"},
+//      {"$P$9saltstriTWMzWKsEeiE7CKOVVU.rS0","abcdefghijklmn"},
+//      {"$P$9saltstriXt7EDPKtkyRVOqcqEW5UU.", "abcdefghijklmno"},
+	{NULL}
+};
+
+static void init(void)
+{
+}
+
+static int valid(char *ciphertext)
+{
+	uint32_t i, count_log2;
+
+	if (strlen(ciphertext) != CIPHERTEXT_LENGTH)
+		return 0;
+	if (strncmp(ciphertext, phpass_prefix, 3) != 0)
+		return 0;
+
+	for (i = 3; i < CIPHERTEXT_LENGTH; i++)
+		if (atoi64[ARCH_INDEX(ciphertext[i])] == 0x7F)
+			return 0;
+
+	count_log2 = atoi64[ARCH_INDEX(ciphertext[3])];
+	if (count_log2 < 7 || count_log2 > 31)
+		return 0;
+
+	return 1;
+};
+
+//code from historical JtR phpass patch
+static void *binary(char *ciphertext)
+{
+	static unsigned char b[BINARY_SIZE];
+	memset(b, 0, BINARY_SIZE);
+	int i, bidx = 0;
+	unsigned sixbits;
+	char *pos = &ciphertext[3 + 1 + 8];
+
+	for (i = 0; i < 5; i++) {
+		sixbits = atoi64[ARCH_INDEX(*pos++)];
+		b[bidx] = sixbits;
+		sixbits = atoi64[ARCH_INDEX(*pos++)];
+		b[bidx++] |= (sixbits << 6);
+		sixbits >>= 2;
+		b[bidx] = sixbits;
+		sixbits = atoi64[ARCH_INDEX(*pos++)];
+		b[bidx++] |= (sixbits << 4);
+		sixbits >>= 4;
+		b[bidx] = sixbits;
+		sixbits = atoi64[ARCH_INDEX(*pos++)];
+		b[bidx++] |= (sixbits << 2);
+	}
+	sixbits = atoi64[ARCH_INDEX(*pos++)];
+	b[bidx] = sixbits;
+	sixbits = atoi64[ARCH_INDEX(*pos++)];
+	b[bidx] |= (sixbits << 6);
+	return (void *) b;
+}
+
+static void *salt(char *ciphertext)
+{
+	static unsigned char salt[SALT_SIZE + 2];
+	memcpy(salt, &ciphertext[4], 8);
+	salt[8] = ciphertext[3];
+	salt[9] = 0;
+	return salt;
+}
+
+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)
+{
+	int i;
+	unsigned char *csalt = salt;
+	for (i = 0; i < 8; i++)
+		currentsalt[i] = csalt[i];
+	loopChar = csalt[8];
+}
+
+static void set_key(char *key, int index)
+{
+	int i, len = strlen(key);
+	inbuffer[address(15, index)] = len;
+	for (i = 0; i < len; i++)
+		inbuffer[address(i, index)] = key[i];
+	inbuffer[address(len, index)] = 0;
+}
+
+static char *get_key(int index)
+{
+	static char r[PLAINTEXT_LENGTH + 1];
+	int i;
+	for (i = 0; i < PLAINTEXT_LENGTH; i++)
+		r[i] = inbuffer[address(i, index)];
+	r[PLAINTEXT_LENGTH] = '\0';
+	return r;
+}
+
+static void crypt_all(int count)
+{
+	char setting[40];
+	strcpy(setting, phpass_prefix);
+	setting[3] = loopChar;
+	int count_log2 = 0;
+	count_log2 = atoi64[ARCH_INDEX(setting[3])];
+	strcpy(setting + 4, currentsalt);
+	mem_init(inbuffer, outbuffer, setting, itoa64, count_log2);
+	gpu_phpass();
+	mem_clear();
+}
+
+static int get_hash_0(int index)
+{
+	return outbuffer[address(0, index)] & 0xf;
+}
+
+static int get_hash_1(int index)
+{
+	return outbuffer[address(0, index)] & 0xff;
+}
+
+static int get_hash_2(int index)
+{
+	return outbuffer[address(0, index)] & 0xfff;
+}
+
+static int get_hash_3(int index)
+{
+	return outbuffer[address(0, index)] & 0xffff;
+}
+
+static int get_hash_4(int index)
+{
+	return outbuffer[address(0, index)] & 0xfffff;
+}
+
+static int cmp_all(void *binary, int count)
+{
+	uint32_t i;
+	uint32_t b = ((uint32_t *) binary)[0];
+	for (i = 0; i < count; i++)
+		if (b == outbuffer[address(0, i)])
+			return 1;
+	return 0;
+}
+
+static int cmp_one(void *binary, int index)
+{
+	int i;
+	uint32_t *t = (uint32_t *) binary;
+	for (i = 0; i < 4; i++)
+		if (t[i] != outbuffer[address(i, index)])
+			return 0;
+	return 1;
+}
+
+static int cmp_exact(char *source, int count)
+{
+	return 1;
+}
+
+struct fmt_main fmt_PHPASSCUDA = {
+	{
+		    FORMAT_LABEL,
+		    FORMAT_NAME,
+		    PHPASS_TYPE,
+		    BENCHMARK_COMMENT,
+		    BENCHMARK_LENGTH,
+		    PLAINTEXT_LENGTH,
+		    BINARY_SIZE,
+		    SALT_SIZE + 1,
+		    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}
+};
