diff -urpN john-1.7.9.orig//src/Makefile john-1.7.9-opencl_2//src/Makefile
--- john-1.7.9.orig//src/Makefile	2011-11-23 00:19:21.000000000 +0000
+++ john-1.7.9-opencl_2//src/Makefile	2011-12-11 16:21:39.682868795 +0000
@@ -20,6 +20,9 @@ OMPFLAGS =
 #OMPFLAGS = -fopenmp -msse2
 # Sun Studio with OpenMP (set the OMP_NUM_THREADS env var at runtime)
 #OMPFLAGS = -xopenmp
+#CUDALIBPATH = /usr/local/cuda/include/
+ATISTREAMSDKROOT = /opt/AMDAPP
+OCLHEADERPATH = -I$(ATISTREAMSDKROOT)/include
 CFLAGS = -c -Wall -O2 -fomit-frame-pointer $(OMPFLAGS)
 # CFLAGS for use on the main john.c file only
 CFLAGS_MAIN = $(CFLAGS)
@@ -38,6 +41,8 @@ JOHN_OBJS = \
 	LM_fmt.o \
 	trip_fmt.o \
 	dummy.o \
+	phpassopencl_fmt.o \
+	cryptmd5opencl_fmt.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 \
@@ -79,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-opencl      Linux, x86-64 with SSE2 and OpenCL (experimental)"
+	@echo "linux-x86-opencl         Linux, x86 32-bit with SSE2 and OpenCL (experimental)"
 	@echo "linux-x86-64-avx         Linux, x86-64 with AVX (2011+ Intel CPUs)"
 	@echo "linux-x86-64-xop         Linux, x86-64 with AVX and XOP (2011+ AMD CPUs)"
 	@echo "linux-x86-64             Linux, x86-64 with SSE2 (most common)"
@@ -131,6 +138,7 @@ default:
 	@echo "sco-x86-any-cc           SCO, x86, cc"
 	@echo "tru64-alpha              Tru64 (Digital UNIX, OSF/1), Alpha"
 	@echo "aix-ppc32                AIX, PowerPC 32-bit"
+	@echo "macosx-x86-64-opencl     Mac OS X 10.5+, Xcode 3.0+, x86-64 with SSE2 and OpenCL (experimental)"
 	@echo "macosx-x86-64            Mac OS X 10.5+, Xcode 3.0+, x86-64 with SSE2 (best)"
 	@echo "macosx-x86-sse2          Mac OS X, x86 with SSE2"
 	@echo "macosx-ppc32-altivec     Mac OS X, PowerPC w/AltiVec (best)"
@@ -153,6 +161,28 @@ default:
 	@echo "beos-x86-any             BeOS, x86"
 	@echo "generic                  Any other Unix-like system with gcc"
 
+linux-x86-64-opencl:
+	$(LN) x86-64.h arch.h
+	$(MAKE) $(PROJ) \
+		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86-64.o" \
+		CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \
+		LDFLAGS="$(LDFLAGS) -lcrypt -L$(ATISTREAMSDKROOT)/lib/x86_64 -lOpenCL"
+
+linux-x86-opencl:
+	$(LN) x86-sse.h arch.h
+	$(MAKE) $(PROJ) \
+		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86.o x86-sse.o" \
+		CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \
+		LDFLAGS="$(LDFLAGS) -lcrypt -L$(ATISTREAMSDKROOT)/lib/x86 -lOpenCL"
+
+macosx-x86-64-opencl:
+	$(LN) x86-64.h arch.h
+	$(MAKE) $(PROJ) \
+		JOHN_OBJS="$(JOHN_OBJS) x86-64.o" \
+		ASFLAGS="$(ASFLAGS) -m64 -DUNDERSCORES -DBSD -DALIGN_LOG" \
+		CFLAGS="$(CFLAGS) -m64" \
+		LDFLAGS="$(LDFLAGS) -m64 -framework OpenCL"
+
 linux-x86-64-avx:
 	$(LN) x86-64.h arch.h
 	$(MAKE) $(PROJ) \
@@ -821,6 +851,12 @@ generic.h:
 		"$(BENCH_MD5_OBJS_DEPEND)" \
 		"$(BENCH_BF_OBJS_DEPEND)"
 
+phpassopencl_fmt.o: phpassopencl_fmt.c
+	$(CC) $(OCLHEADERPATH) $(CFLAGS) phpassopencl_fmt.c
+
+cryptmd5opencl_fmt.o: cryptmd5opencl_fmt.c
+	$(CC) $(OCLHEADERPATH) $(CFLAGS) cryptmd5opencl_fmt.c
+
 bench: $(BENCH_OBJS)
 	$(LD) $(BENCH_OBJS) $(LDFLAGS) -o bench
 
@@ -903,8 +939,11 @@ 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
 
+nice:
+	 indent -kr -i8 -nlp -nbbo -l79 -lc79 phpassopencl_fmt.c cryptmd5opencl_fmt.c phpass.cl cryptmd5.cl common_opencl.h
+
 include Makefile.dep
diff -urpN john-1.7.9.orig//src/common_opencl.h john-1.7.9-opencl_2//src/common_opencl.h
--- john-1.7.9.orig//src/common_opencl.h	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.9-opencl_2//src/common_opencl.h	2011-12-18 04:17:42.317019084 +0000
@@ -0,0 +1,158 @@
+/*
+* 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 _COMMON_OPENCL_H
+#define _COMMON_OPENCL_H
+#include "path.h"
+static char *get_error_name(cl_int cl_error);
+
+static void HandleError(cl_int cl_error, const char *file, int line)
+{
+	if (cl_error != CL_SUCCESS) {
+		fprintf(stderr, "Error (%s) in file (%s) at line (%d)\n",
+		    get_error_name(cl_error), file, line);
+		exit(EXIT_FAILURE);
+	}
+}
+
+#define HANDLE_ERROR(cl_error) (HandleError(cl_error,__FILE__,__LINE__))
+
+#define SOURCE_BUFSIZE 1024*64
+
+static char kernel_source[SOURCE_BUFSIZE];
+static cl_context context;
+static cl_command_queue cq;
+static cl_program prog;
+static size_t src_size = 0;
+
+
+static void read_kernel_source(char *src, char *kernel_filename)
+{
+	static const char path_prefix[] = "$JOHN/../src/";
+	static char fullname[32];
+	memcpy(fullname, path_prefix, sizeof(path_prefix) - 1);
+	memcpy(fullname + sizeof(path_prefix) - 1, kernel_filename,
+	    strlen(kernel_filename) + 1);
+	char *kernel_path = path_expand(fullname);
+	FILE *f = fopen(kernel_path, "r");
+	if (f != NULL) {
+		src_size = fread(src, 1, SOURCE_BUFSIZE, f);
+		fclose(f);
+	}
+}
+
+static void opencl_init(char *kernel_name)
+{
+	cl_int cl_error;
+	cl_platform_id platform;
+	cl_device_id device;
+	cl_uint platforms, devices;
+	static char log[64];
+
+	///Read kernel source from file
+	read_kernel_source(kernel_source, kernel_name);
+
+	///Find GPU's
+	HANDLE_ERROR(clGetPlatformIDs(1, &platform, &platforms));
+	HANDLE_ERROR(clGetDeviceIDs
+	    (platform, CL_DEVICE_TYPE_GPU, 1, &device, &devices));
+	fprintf(stdout, "Found %d device(s)\n", devices);
+	cl_context_properties properties[] = {
+		CL_CONTEXT_PLATFORM, (cl_context_properties) platform,
+		0
+	};
+	HANDLE_ERROR(clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(log), log,
+		NULL));
+	fprintf(stdout, "\nOpenCL Device: %s\n", log);
+
+	///Setup context
+	context =
+	    clCreateContext(properties, 1, &device, NULL, NULL, &cl_error);
+	cq = clCreateCommandQueue(context, device, 0, &cl_error);
+	///Build kernel
+	const char *srcptr[] = { kernel_source };
+	prog =
+	    clCreateProgramWithSource(context, 1, srcptr, &src_size,
+	    &cl_error);
+	if (cl_error != CL_SUCCESS) {
+		fprintf(stderr, "Error while creating program [%d]\n",
+		    cl_error);
+		exit(1);
+	}
+	cl_error = clBuildProgram(prog, 0, NULL, "-DFLAG1_ENABLED -cl-opt-disable -cl-strict-aliasing -w", NULL, NULL);
+
+	///Report build errors and warnings
+	cl_build_status build_status;
+	HANDLE_ERROR(clGetProgramBuildInfo(prog, device,
+		CL_PROGRAM_BUILD_STATUS,
+		sizeof(cl_build_status), &build_status, NULL));
+
+	char *build_log;
+	size_t ret_val_size;
+	HANDLE_ERROR(clGetProgramBuildInfo(prog, device, CL_PROGRAM_BUILD_LOG,
+		0, NULL, &ret_val_size));
+
+	build_log = malloc(ret_val_size + 1);
+	HANDLE_ERROR(clGetProgramBuildInfo(prog, device, CL_PROGRAM_BUILD_LOG,
+		ret_val_size, build_log, NULL));
+	build_log[ret_val_size] = '\0';
+	if (strlen(build_log) > 0) {
+		fprintf(stderr, "BUILD LOG:\n %s", build_log);
+	}
+	free(build_log);
+	if (cl_error != CL_SUCCESS) {
+		fprintf(stderr, "Error while building program [%d]\n",
+		    cl_error);
+		exit(1);
+	}
+}
+
+
+static char *get_error_name(cl_int cl_error)
+{
+	static char *err_1[] =
+	    { "CL_SUCCESS", "CL_DEVICE_NOT_FOUND", "CL_DEVICE_NOT_AVAILABLE",
+		"CL_COMPILER_NOT_AVAILABLE",
+		"CL_MEM_OBJECT_ALLOCATION_FAILURE", "CL_OUT_OF_RESOURCES",
+		"CL_OUT_OF_HOST_MEMORY",
+		"CL_PROFILING_INFO_NOT_AVAILABLE", "CL_MEM_COPY_OVERLAP",
+		"CL_IMAGE_FORMAT_MISMATCH",
+		"CL_IMAGE_FORMAT_NOT_SUPPORTED", "CL_BUILD_PROGRAM_FAILURE",
+		"CL_MAP_FAILURE"
+	};
+	static char *err_invalid[] = {
+		"CL_INVALID_VALUE", "CL_INVALID_DEVICE_TYPE",
+		"CL_INVALID_PLATFORM", "CL_INVALID_DEVICE",
+		"CL_INVALID_CONTEXT", "CL_INVALID_QUEUE_PROPERTIES",
+		"CL_INVALID_COMMAND_QUEUE", "CL_INVALID_HOST_PTR",
+		"CL_INVALID_MEM_OBJECT", "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR",
+		"CL_INVALID_IMAGE_SIZE", "CL_INVALID_SAMPLER",
+		"CL_INVALID_BINARY", "CL_INVALID_BUILD_OPTIONS",
+		"CL_INVALID_PROGRAM", "CL_INVALID_PROGRAM_EXECUTABLE",
+		"CL_INVALID_KERNEL_NAME", "CL_INVALID_KERNEL_DEFINITION",
+		"CL_INVALID_KERNEL", "CL_INVALID_ARG_INDEX",
+		"CL_INVALID_ARG_VALUE", "CL_INVALID_ARG_SIZE",
+		"CL_INVALID_KERNEL_ARGS", "CL_INVALID_WORK_DIMENSION",
+		"CL_INVALID_WORK_GROUP_SIZE", "CL_INVALID_WORK_ITEM_SIZE",
+		"CL_INVALID_GLOBAL_OFFSET", "CL_INVALID_EVENT_WAIT_LIST",
+		"CL_INVALID_EVENT", "CL_INVALID_OPERATION",
+		"CL_INVALID_GL_OBJECT", "CL_INVALID_BUFFER_SIZE",
+		"CL_INVALID_MIP_LEVEL", "CL_INVALID_GLOBAL_WORK_SIZE"
+	};
+
+	if (cl_error <= 0 && cl_error >= -12) {
+		cl_error = -cl_error;
+		return err_1[cl_error];
+	}
+	if (cl_error <= -30 && cl_error >= -63) {
+		cl_error = -cl_error;
+		return err_invalid[cl_error - 30];
+	}
+
+	return "UNKNOWN ERROR :(";
+}
+
+
+#endif
diff -urpN john-1.7.9.orig//src/cryptmd5.cl john-1.7.9-opencl_2//src/cryptmd5.cl
--- john-1.7.9.orig//src/cryptmd5.cl	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.9-opencl_2//src/cryptmd5.cl	2011-12-16 03:55:05.817178408 +0000
@@ -0,0 +1,303 @@
+/*
+* 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.
+*/
+#define uint32_t	unsigned int
+#define uint8_t		unsigned char
+
+#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 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
+
+typedef struct {
+	uint8_t saltlen;
+	uint8_t salt[8];
+	uint8_t prefix;		/** 'a' when $apr1$ or '1' when $1$ **/
+} crypt_md5_salt;
+
+typedef struct {
+	uint8_t length;
+	uint8_t v[15];
+} crypt_md5_password;
+
+typedef struct {
+	uint32_t v[4];		/** 128 bits **/
+} crypt_md5_hash;
+
+typedef struct {
+#define ctx_buffsize 64
+	uint8_t buffer[ctx_buffsize];
+	uint32_t buflen;
+	uint32_t len;
+	uint32_t A, B, C, D;
+} md5_ctx;
+
+
+__constant uint8_t cl_md5_salt_prefix[] = "$1$";
+__constant uint8_t cl_apr1_salt_prefix[] = "$apr1$";
+
+void ctx_update_global(__private md5_ctx * ctx, __global uint8_t * string,
+    size_t len)
+{
+	uint8_t *dest = &ctx->buffer[ctx->buflen];
+	__global uint8_t *src = string;
+	ctx->buflen += len;
+	int i = len;
+	for (i = 0; i < len; i++)
+		dest[i] = src[i];
+}
+
+void ctx_update_private(__private md5_ctx * ctx, __private uint8_t * string,
+    size_t len)
+{
+	uint8_t *dest = &ctx->buffer[ctx->buflen];
+	__private uint8_t *src = string;
+	ctx->buflen += len;
+	int i = len;
+	for (i = 0; i < len; i++)
+		dest[i] = src[i];
+}
+
+void ctx_update_prefix(__private md5_ctx * ctx, uint8_t prefix)
+{
+	uint8_t i, *dest = &ctx->buffer[ctx->buflen];
+	if (prefix == '1') {
+		ctx->buflen += 3;
+		for (i = 0; i < 3; i++)
+			dest[i] = cl_md5_salt_prefix[i];
+	} else {
+		ctx->buflen += 6;
+		for (i = 0; i < 6; i++)
+			dest[i] = cl_apr1_salt_prefix[i];
+	}
+}
+
+
+void init_ctx(__private md5_ctx * ctx)
+{
+	int i = ctx_buffsize / sizeof(uint32_t);
+	uint32_t *buf = (uint32_t *) ctx->buffer;
+	while (i--)
+		*buf++ = 0;
+	ctx->buflen = 0;
+	ctx->len = 0;
+}
+
+void md5_block(__private md5_ctx * ctx, uint32_t blocks, size_t len)
+{
+	uint32_t a = 0x67452301;
+	uint32_t b = 0xefcdab89;
+	uint32_t c = 0x98badcfe;
+	uint32_t d = 0x10325476;
+
+	ctx->len += len;
+	len <<= 3;
+	__private uint32_t *x = (uint32_t *) & ctx->buffer[0];
+
+	{
+		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 */
+		FF(b, c, d, a, 0, 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 */
+		GG(c, d, a, b, 0, 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 */
+		HH(c, d, a, b, 0, 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 */
+		II(d, a, b, c, 0, 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 */
+	}
+	ctx->A = a + 0x67452301;
+	ctx->B = b + 0xefcdab89;
+	ctx->C = c + 0x98badcfe;
+	ctx->D = d + 0x10325476;
+}
+
+
+void md5_digest(__private md5_ctx * ctx, __private uint32_t * result)
+{
+	uint32_t len = ctx->buflen, blocks = 1;
+	uint32_t *x = (uint32_t *) ctx->buffer;
+	uint32_t i = len % 64;
+	x[i / 4] |= (((uint32_t) 0x80) << ((i & 0x3) << 3));
+
+	md5_block(ctx, blocks, len);
+
+	result[0] = ctx->A;
+	result[1] = ctx->B;
+	result[2] = ctx->C;
+	result[3] = ctx->D;
+}
+
+__kernel void cryptmd5
+    (__global const crypt_md5_password * inbuffer,
+    __global crypt_md5_hash * outbuffer,
+    __global const crypt_md5_salt * hsalt) {
+	uint32_t idx = get_global_id(0);
+	uint32_t i;
+	__global const uint8_t *pass = inbuffer[idx].v;
+	__global uint32_t *tresult = outbuffer[idx].v;
+
+	__private uint32_t alt_result[4];
+	uint8_t pass_len = inbuffer[idx].length;
+	uint8_t salt_len = hsalt->saltlen;
+	const __global uint8_t *salt = hsalt->salt;
+
+	__private md5_ctx ctx, alt_ctx;
+	init_ctx(&ctx);
+	init_ctx(&alt_ctx);
+
+	ctx_update_global(&ctx, (__global uint8_t *) pass, pass_len);
+	ctx_update_prefix(&ctx, hsalt->prefix);
+	ctx_update_global(&ctx, (__global uint8_t *) salt, salt_len);
+
+	ctx_update_global(&alt_ctx, (__global uint8_t *) pass, pass_len);
+	ctx_update_global(&alt_ctx, (__global uint8_t *) salt, salt_len);
+	ctx_update_global(&alt_ctx, (__global uint8_t *) pass, pass_len);
+	md5_digest(&alt_ctx, alt_result);
+
+	for (i = pass_len; i > 16; i -= 16)
+		ctx_update_private(&ctx, (uint8_t *) alt_result, 16);
+	ctx_update_private(&ctx, (uint8_t *) alt_result, i);
+
+
+	*alt_result = 0;
+
+	for (i = pass_len; i > 0; i >>= 1)
+		if ((i & 1) != 0)
+			ctx.buffer[ctx.buflen++] = ((char *) alt_result)[0];
+		else
+			ctx.buffer[ctx.buflen++] = pass[0];
+
+	md5_digest(&ctx, alt_result);
+
+
+
+	for (i = 0; i < 1000; i++) {
+		init_ctx(&ctx);
+
+		if ((i & 1) != 0)
+			ctx_update_global(&ctx, (__global uint8_t *) pass,
+			    pass_len);
+		else
+			ctx_update_private(&ctx, (uint8_t *) alt_result, 16);
+
+		if (i % 3 != 0)
+			ctx_update_global(&ctx, (__global uint8_t *) salt,
+			    salt_len);
+
+		if (i % 7 != 0)
+			ctx_update_global(&ctx, (__global uint8_t *) pass,
+			    pass_len);
+
+		if ((i & 1) != 0)
+			ctx_update_private(&ctx, (uint8_t *) alt_result, 16);
+		else
+			ctx_update_global(&ctx, (__global uint8_t *) pass,
+			    pass_len);
+		md5_digest(&ctx, alt_result);
+
+	}
+	tresult[0] = ctx.A;
+	tresult[1] = ctx.B;
+	tresult[2] = ctx.C;
+	tresult[3] = ctx.D;
+}
diff -urpN john-1.7.9.orig//src/cryptmd5opencl_fmt.c john-1.7.9-opencl_2//src/cryptmd5opencl_fmt.c
--- john-1.7.9.orig//src/cryptmd5opencl_fmt.c	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.9-opencl_2//src/cryptmd5opencl_fmt.c	2012-01-15 02:46:48.631919627 +0000
@@ -0,0 +1,417 @@
+/*
+* 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 "path.h"
+
+#ifdef __APPLE__
+#include <OpenCL/opencl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+#include "common_opencl.h"
+
+#define uint32_t unsigned int
+#define uint8_t unsigned char
+
+#define BLOCKS			18
+#define THREADS 		256
+#define KEYS_PER_CRYPT 		BLOCKS*THREADS
+#define PLAINTEXT_LENGTH	15
+
+#define MIN(a,b) 		((a)<(b)?(a):(b))
+#define MAX(a,b) 		((a)>(b)?(a):(b))
+
+#define FORMAT_LABEL		"cryptmd5opencl"
+#define FORMAT_NAME		"CRYPTMD5OPENCL"
+#define KERNEL_NAME		"cryptmd5"
+
+#define CRYPT_TYPE		"MD5-based CRYPT"
+
+#define BENCHMARK_COMMENT	""
+#define BENCHMARK_LENGTH	-1
+
+#define BINARY_SIZE		16
+#define SALT_SIZE		(8+1)					/** salt + prefix id **/
+#define MIN_KEYS_PER_CRYPT	KEYS_PER_CRYPT
+#define MAX_KEYS_PER_CRYPT	KEYS_PER_CRYPT
+
+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[ctx_buffsize];
+	uint32_t buflen;
+	uint32_t len;
+	uint32_t A, B, C, D;
+} md5_ctx;
+
+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 **/
+
+static const char md5_salt_prefix[] = "$1$";
+static const char apr1_salt_prefix[] = "$apr1$";
+
+
+//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)
+{
+	opencl_init("cryptmd5.cl");
+}
+
+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) \
+  {\
+      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;\
+  }
+
+	_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)
+{
+	static char ret[PLAINTEXT_LENGTH + 1];
+	memcpy(ret, inbuffer[index].v, PLAINTEXT_LENGTH);
+	ret[inbuffer[index].length] = '\0';
+	return ret;
+}
+
+static void crypt_all(int count)
+{
+	size_t insize = sizeof(crypt_md5_password) * KEYS_PER_CRYPT;
+	size_t outsize = sizeof(crypt_md5_hash) * KEYS_PER_CRYPT;
+	size_t saltsize = sizeof(crypt_md5_salt);
+
+	///Alocate memory on the GPU
+	cl_int cl_error;
+	cl_mem mem_in, mem_out, mem_salt;
+	mem_in =
+	    clCreateBuffer(context, CL_MEM_READ_ONLY, insize, NULL, &cl_error);
+	mem_salt =
+	    clCreateBuffer(context, CL_MEM_READ_ONLY, saltsize, NULL,
+	    &cl_error);
+	mem_out =
+	    clCreateBuffer(context, CL_MEM_WRITE_ONLY, outsize, NULL,
+	    &cl_error);
+
+	///Assign kernel parameters 
+	cl_kernel kernel = clCreateKernel(prog, KERNEL_NAME, &cl_error);
+	clSetKernelArg(kernel, 0, sizeof(mem_in), &mem_in);
+	clSetKernelArg(kernel, 1, sizeof(mem_out), &mem_out);
+	clSetKernelArg(kernel, 2, sizeof(mem_salt), &mem_salt);
+
+	///Copy data to GPU memory
+	HANDLE_ERROR(clEnqueueWriteBuffer
+	    (cq, mem_in, CL_FALSE, 0, insize, inbuffer, 0, NULL, NULL));
+	HANDLE_ERROR(clEnqueueWriteBuffer
+	    (cq, mem_salt, CL_FALSE, 0, saltsize, &host_salt, 0, NULL, NULL));
+
+	///Run kernel
+	size_t worksize = KEYS_PER_CRYPT;
+	size_t localworksize = THREADS;
+	HANDLE_ERROR(clEnqueueNDRangeKernel
+	    (cq, kernel, 1, NULL, &worksize, &localworksize, 0, NULL, NULL));
+	HANDLE_ERROR(clEnqueueReadBuffer
+	    (cq, mem_out, CL_FALSE, 0, outsize, outbuffer, 0, NULL, NULL));
+
+	///Await completion of all the above
+	HANDLE_ERROR(clFinish(cq));
+
+	/// Cleanup
+	HANDLE_ERROR(clReleaseKernel(kernel));
+	HANDLE_ERROR(clReleaseMemObject(mem_in));
+	HANDLE_ERROR(clReleaseMemObject(mem_salt));
+	HANDLE_ERROR(clReleaseMemObject(mem_out));
+
+}
+
+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_CRYPTMD5OPENCL = {
+	{
+		    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.9.orig//src/john.c john-1.7.9-opencl_2//src/john.c
--- john-1.7.9.orig//src/john.c	2011-11-23 02:46:59.000000000 +0000
+++ john-1.7.9-opencl_2//src/john.c	2011-12-10 08:39:08.620216589 +0000
@@ -44,6 +44,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_PHPASSOPENCL,fmt_CRYPTMD5OPENCL;
 #ifdef HAVE_CRYPT
 extern struct fmt_main fmt_crypt;
 #endif
@@ -77,6 +78,8 @@ 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_PHPASSOPENCL);
+	john_register_one(&fmt_CRYPTMD5OPENCL);
 #ifdef HAVE_CRYPT
 	john_register_one(&fmt_crypt);
 #endif
diff -urpN john-1.7.9.orig//src/phpass.cl john-1.7.9-opencl_2//src/phpass.cl
--- john-1.7.9.orig//src/phpass.cl	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.9-opencl_2//src/phpass.cl	2011-12-17 12:44:02.302637632 +0000
@@ -0,0 +1,336 @@
+/*
+* 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.
+*/
+#define PLAINTEXT_LENGTH	15
+typedef struct {
+	unsigned char v[PLAINTEXT_LENGTH];
+	unsigned char length;
+} phpass_password;
+
+typedef struct {
+	unsigned int v[4];
+} phpass_hash;
+
+
+#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(a, b, c, d, x, s, ac) \
+  {(a) += F ((b), (c), (d)) + (x) + (uint2)(ac); \
+   (a) = ROTATE_LEFT ((a), (s)); \
+   (a) += (b); \
+  }
+#define GG(a, b, c, d, x, s, ac) \
+  {(a) += G ((b), (c), (d)) + (x) + (uint2)(ac); \
+   (a) = ROTATE_LEFT ((a), (s)); \
+   (a) += (b); \
+  }
+#define HH(a, b, c, d, x, s, ac) \
+  {(a) += H ((b), (c), (d)) + (x) + (uint2)(ac); \
+   (a) = ROTATE_LEFT ((a), (s)); \
+   (a) += (b); \
+  }
+#define II(a, b, c, d, x, s, ac) \
+  {(a) += I ((b), (c), (d)) + (x) + (uint2)(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 SALT_SIZE		8
+
+#define AC1			0xd76aa477
+#define AC2pCd			0xf8fa0bcc
+#define AC3pCc			0xbcdb4dd9
+#define AC4pCb			0xb18b7a77
+#define MASK1			0x77777777
+
+
+inline void md5(uchar2 len, __private uint2 * internal_ret,
+    __private uint2 * x)
+{
+	uint t;
+#define K(q)\
+	t=x[len.s##q/4].s##q;\
+	t|= (((uint) 0x80) << ((len.s##q & 0x3) << 3));\
+	x[len.s##q/4].s##q=t;
+	K(0);
+	K(1);
+#undef K
+	uint2 x14 = (uint2) (len.s0 << 3, len.s1 << 3);
+
+	uint2 a = 0x67452301;
+	uint2 b = 0xefcdab89;
+	uint2 c = 0x98badcfe;
+	uint2 d = 0x10325476;
+
+	FF(a, b, c, d, x[0], S11, 0xd76aa478);
+	FF(d, a, b, c, x[1], S12, 0xe8c7b756);
+	FF(c, d, a, b, x[2], S13, 0x242070db);
+	FF(b, c, d, a, x[3], S14, 0xc1bdceee);
+	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, 0, S11, 0x698098d8);
+	FF(d, a, b, c, 0, S12, 0x8b44f7af);
+	FF(c, d, a, b, 0, S13, 0xffff5bb1);
+	FF(b, c, d, a, 0, S14, 0x895cd7be);
+	FF(a, b, c, d, 0, S11, 0x6b901122);
+	FF(d, a, b, c, 0, 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, 0, 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, 0, S22, 0x2441453);
+	GG(c, d, a, b, 0, S23, 0xd8a1e681);
+	GG(b, c, d, a, x[4], S24, 0xe7d3fbc8);
+	GG(a, b, c, d, 0, S21, 0x21e1cde6);
+	GG(d, a, b, c, x14, S22, 0xc33707d6);
+	GG(c, d, a, b, x[3], S23, 0xf4d50d87);
+	GG(b, c, d, a, 0, S24, 0x455a14ed);
+	GG(a, b, c, d, 0, 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, 0, S24, 0x8d2a4c8a);
+
+	HH(a, b, c, d, x[5], S31, 0xfffa3942);
+	HH(d, a, b, c, 0, S32, 0x8771f681);
+	HH(c, d, a, b, 0, 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, 0, S34, 0xbebfbc70);
+	HH(a, b, c, d, 0, 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, 0, S31, 0xd9d4d039);
+	HH(d, a, b, c, 0, 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, 0, S41, 0x655b59c3);
+	II(d, a, b, c, x[3], S42, 0x8f0ccc92);
+	II(c, d, a, b, 0, S43, 0xffeff47d);
+	II(b, c, d, a, x[1], S44, 0x85845dd1);
+	II(a, b, c, d, 0, S41, 0x6fa87e4f);
+	II(d, a, b, c, 0, S42, 0xfe2ce6e0);
+	II(c, d, a, b, x[6], S43, 0xa3014314);
+	II(b, c, d, a, 0, S44, 0x4e0811a1);
+	II(a, b, c, d, x[4], S41, 0xf7537e82);
+	II(d, a, b, c, 0, S42, 0xbd3af235);
+	II(c, d, a, b, x[2], S43, 0x2ad7d2bb);
+	II(b, c, d, a, 0, S44, 0xeb86d391);
+
+	internal_ret[0] = a + (uint2) (0x67452301);
+	internal_ret[1] = b + (uint2) (0xefcdab89);
+	internal_ret[2] = c + (uint2) (0x98badcfe);
+	internal_ret[3] = d + (uint2) (0x10325476);
+}
+
+inline void clear_ctx(__private uint2 * x)
+{
+	int i;
+	uint2 zero = (uint2) (0, 0);
+	for (i = 0; i < 8; i++)
+		*x++ = zero;
+}
+inline void clean_ctx(__private uint *x){
+	int i;
+	for(i=0;i<8;i++);
+		*x++=0;
+}
+
+
+__kernel void phpass
+    (__global const phpass_password * data, __global phpass_hash * data_out,
+    __global const char *setting) {
+	uint2 x[8];
+	clear_ctx(x);
+
+	uint idx = get_global_id(0);
+
+	__global const uchar *password0 = data[idx * 2].v;
+	__global const uchar *password1 = data[idx * 2 + 1].v;
+
+	uchar2 length;		//password length
+	length.x = data[idx * 2 + 0].length;
+	length.y = data[idx * 2 + 1].length;
+	
+	uint count, i = 0;
+	
+	uint sx[8];
+	__private uchar *buff = (uchar *) sx;
+ 	#define K(q)\
+		clean_ctx(sx);\
+		for (i = 0; i < 8; i++)\
+			buff[i] = setting[i];\
+		for (i = 8; i < 8 + length.s##q; i++)\
+			buff[i] = password##q[i - 8];\
+		for ( i = 0; i < 8; i++)\
+			x[i].s##q=sx[i];
+		K(0);
+		K(1);
+	#undef K
+
+	
+	uchar2 length2 = length + (uchar2) (8);
+	md5(length2, x, x);
+	#define K(q)\
+		clean_ctx(sx);\
+		for(i=0;i<length.s##q;i++)\
+			buff[i]=password##q[i];\
+		for(i=0;i<4;i++)\
+			x[i+4].s##q=sx[i];
+	K(0);
+	K(1);
+	#undef K
+
+	count = 1 << setting[SALT_SIZE + 3];
+
+	uchar2 length3 = length+(uchar2)(16);
+
+	uint2 a, b, c, d, x0, x1, x2, x3, x4, x5, x6, x7;
+	uint2 x14 = (uint2)(length3.s0 << 3,length3.s1 << 3);
+
+	uint t;
+#define K(q)\
+	t=x[length3.s##q/4].s##q;\
+	t|= (((uint) 0x80) << ((length3.s##q & 0x3) << 3));\
+	x[length3.s##q/4].s##q=t;
+	K(0);
+	K(1);
+#undef K
+
+	x0 = x[0];
+	x1 = x[1];
+	x2 = x[2];
+	x3 = x[3];
+	x4 = x[4];
+	x5 = x[5];
+	x6 = x[6];
+	x7 = x[7];
+
+	do {
+	   b = 0xefcdab89;
+	   c = 0x98badcfe;
+	   d = 0x10325476;
+
+	   FF(a, b, c, d, x0, S11, 0xd76aa478);
+	   FF(d, a, b, c, x1, S12, 0xe8c7b756);
+	   FF(c, d, a, b, x2, S13, 0x242070db);
+	   FF(b, c, d, a, x3, S14, 0xc1bdceee);
+	   FF(a, b, c, d, x4, S11, 0xf57c0faf);
+	   FF(d, a, b, c, x5, S12, 0x4787c62a);
+	   FF(c, d, a, b, x6, S13, 0xa8304613);
+	   FF(b, c, d, a, x7, S14, 0xfd469501);
+	   FF(a, b, c, d, 0, S11, 0x698098d8);
+	   FF(d, a, b, c, 0, S12, 0x8b44f7af);
+	   FF(c, d, a, b, 0, S13, 0xffff5bb1);
+	   FF(b, c, d, a, 0, S14, 0x895cd7be);
+	   FF(a, b, c, d, 0, S11, 0x6b901122);
+	   FF(d, a, b, c, 0, S12, 0xfd987193);
+	   FF(c, d, a, b, x14, S13, 0xa679438e);
+	   FF(b, c, d, a, 0, S14, 0x49b40821);
+
+	   GG(a, b, c, d, x1, S21, 0xf61e2562);
+	   GG(d, a, b, c, x6, S22, 0xc040b340);
+	   GG(c, d, a, b, 0, S23, 0x265e5a51);
+	   GG(b, c, d, a, x0, S24, 0xe9b6c7aa);
+	   GG(a, b, c, d, x5, S21, 0xd62f105d);
+	   GG(d, a, b, c, 0, S22, 0x2441453);
+	   GG(c, d, a, b, 0, S23, 0xd8a1e681);
+	   GG(b, c, d, a, x4, S24, 0xe7d3fbc8);
+	   GG(a, b, c, d, 0, S21, 0x21e1cde6);
+	   GG(d, a, b, c, x14, S22, 0xc33707d6);
+	   GG(c, d, a, b, x3, S23, 0xf4d50d87);
+	   GG(b, c, d, a, 0, S24, 0x455a14ed);
+	   GG(a, b, c, d, 0, S21, 0xa9e3e905);
+	   GG(d, a, b, c, x2, S22, 0xfcefa3f8);
+	   GG(c, d, a, b, x7, S23, 0x676f02d9);
+	   GG(b, c, d, a, 0, S24, 0x8d2a4c8a);
+
+	   HH(a, b, c, d, x5, S31, 0xfffa3942);
+	   HH(d, a, b, c, 0, S32, 0x8771f681);
+	   HH(c, d, a, b, 0, S33, 0x6d9d6122);
+	   HH(b, c, d, a, x14, S34, 0xfde5380c);
+	   HH(a, b, c, d, x1, S31, 0xa4beea44);
+	   HH(d, a, b, c, x4, S32, 0x4bdecfa9);
+	   HH(c, d, a, b, x7, S33, 0xf6bb4b60);
+	   HH(b, c, d, a, 0, S34, 0xbebfbc70);
+	   HH(a, b, c, d, 0, S31, 0x289b7ec6);
+	   HH(d, a, b, c, x0, S32, 0xeaa127fa);
+	   HH(c, d, a, b, x3, S33, 0xd4ef3085);
+	   HH(b, c, d, a, x6, S34, 0x4881d05);
+	   HH(a, b, c, d, 0, S31, 0xd9d4d039);
+	   HH(d, a, b, c, 0, S32, 0xe6db99e5);
+	   HH(c, d, a, b, 0, S33, 0x1fa27cf8);
+	   HH(b, c, d, a, x2, S34, 0xc4ac5665);
+
+	   II(a, b, c, d, x0, S41, 0xf4292244);
+	   II(d, a, b, c, x7, S42, 0x432aff97);
+	   II(c, d, a, b, x14, S43, 0xab9423a7);
+	   II(b, c, d, a, x5, S44, 0xfc93a039);
+	   II(a, b, c, d, 0, S41, 0x655b59c3);
+	   II(d, a, b, c, x3, S42, 0x8f0ccc92);
+	   II(c, d, a, b, 0, S43, 0xffeff47d);
+	   II(b, c, d, a, x1, S44, 0x85845dd1);
+	   II(a, b, c, d, 0, S41, 0x6fa87e4f);
+	   II(d, a, b, c, 0, S42, 0xfe2ce6e0);
+	   II(c, d, a, b, x6, S43, 0xa3014314);
+	   II(b, c, d, a, 0, S44, 0x4e0811a1);
+	   II(a, b, c, d, x4, S41, 0xf7537e82);
+	   II(d, a, b, c, 0, S42, 0xbd3af235);
+	   II(c, d, a, b, x2, S43, 0x2ad7d2bb);
+	   II(b, c, d, a, 0, S44, 0xeb86d391);
+
+	   x0 = a + (uint2)(0x67452301);
+	   x1 = b + (uint2)(0xefcdab89);
+	   x2 = c + (uint2)(0x98badcfe);
+	   x3 = d + (uint2)(0x10325476);
+
+	   } while (--count);
+	 
+	data_out[idx * 2 + 0].v[0] = x0.s0;
+	data_out[idx * 2 + 0].v[1] = x1.s0;
+	data_out[idx * 2 + 0].v[2] = x2.s0;
+	data_out[idx * 2 + 0].v[3] = x3.s0;
+
+	data_out[idx * 2 + 1].v[0] = x0.s1;
+	data_out[idx * 2 + 1].v[1] = x1.s1;
+	data_out[idx * 2 + 1].v[2] = x2.s1;
+	data_out[idx * 2 + 1].v[3] = x3.s1;
+}
diff -urpN john-1.7.9.orig//src/phpass2.cl john-1.7.9-opencl_2//src/phpass2.cl
--- john-1.7.9.orig//src/phpass2.cl	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.9-opencl_2//src/phpass2.cl	2011-12-18 03:54:54.481065162 +0000
@@ -0,0 +1,388 @@
+/*
+* 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.
+*/
+#define PLAINTEXT_LENGTH	15
+typedef struct {
+	unsigned char v[PLAINTEXT_LENGTH];
+	unsigned char length;
+} phpass_password;
+
+typedef struct {
+	unsigned int v[4];
+} phpass_hash;
+
+
+#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(a, b, c, d, x, s, ac) \
+  {(a) += F ((b), (c), (d)) + (x) + (uint4)(ac); \
+   (a) = ROTATE_LEFT ((a), (s)); \
+   (a) += (b); \
+  }
+#define GG(a, b, c, d, x, s, ac) \
+  {(a) += G ((b), (c), (d)) + (x) + (uint4)(ac); \
+   (a) = ROTATE_LEFT ((a), (s)); \
+   (a) += (b); \
+  }
+#define HH(a, b, c, d, x, s, ac) \
+  {(a) += H ((b), (c), (d)) + (x) + (uint4)(ac); \
+   (a) = ROTATE_LEFT ((a), (s)); \
+   (a) += (b); \
+  }
+#define II(a, b, c, d, x, s, ac) \
+  {(a) += I ((b), (c), (d)) + (x) + (uint4)(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 SALT_SIZE		8
+
+#define AC1			0xd76aa477
+#define AC2pCd			0xf8fa0bcc
+#define AC3pCc			0xbcdb4dd9
+#define AC4pCb			0xb18b7a77
+#define MASK1			0x77777777
+
+
+inline void cuda_md5(uchar4 len, __private uint4 * internal_ret,
+    __private uint4 * x)
+{
+uint t;
+#define K(q)\
+	t=x[len.s##q/4].s##q;\
+	t|= (((uint) 0x80) << ((len.s##q & 0x3) << 3));\
+	x[len.s##q/4].s##q=t;
+	K(0);
+	K(1);
+	K(2);
+	K(3);
+#undef K
+	//x[len.s0 / 4].s0 |= (((uint) 0x80) << ((len.s0 & 0x3) << 3));
+	//x[len.s1 / 4].s1 |= (((uint) 0x80) << ((len.s1 & 0x3) << 3));
+	//x[len.s2 / 4].s2 |= (((uint) 0x80) << ((len.s2 & 0x3) << 3));
+	//x[len.s3 / 4].s3 |= (((uint) 0x80) << ((len.s3 & 0x3) << 3));
+//
+	uint4 x14 = (uint4) (len.s0 << 3, len.s1 << 3, len.s2 << 3, len.s3 << 3);
+
+	uint4 a = (uint4) (0x67452301);
+	uint4 b = (uint4) (0xefcdab89);
+	uint4 c = (uint4) (0x98badcfe);
+	uint4 d = (uint4) (0x10325476);
+
+
+
+	FF(a, b, c, d, x[0], S11, 0xd76aa478);
+    FF(d, a, b, c, x[1], S12, 0xe8c7b756);
+    FF(c, d, a, b, x[2], S13, 0x242070db);
+    FF(b, c, d, a, x[3], S14, 0xc1bdceee);
+	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, 0, S11, 0x698098d8);
+	FF(d, a, b, c, 0, S12, 0x8b44f7af);
+	FF(c, d, a, b, 0, S13, 0xffff5bb1);
+	FF(b, c, d, a, 0, S14, 0x895cd7be);
+	FF(a, b, c, d, 0, S11, 0x6b901122);
+	FF(d, a, b, c, 0, 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, 0, 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, 0, S22, 0x2441453);
+	GG(c, d, a, b, 0, S23, 0xd8a1e681);
+	GG(b, c, d, a, x[4], S24, 0xe7d3fbc8);
+	GG(a, b, c, d, 0, S21, 0x21e1cde6);
+	GG(d, a, b, c, x14, S22, 0xc33707d6);
+	GG(c, d, a, b, x[3], S23, 0xf4d50d87);
+	GG(b, c, d, a, 0, S24, 0x455a14ed);
+	GG(a, b, c, d, 0, 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, 0, S24, 0x8d2a4c8a);
+
+	HH(a, b, c, d, x[5], S31, 0xfffa3942);
+	HH(d, a, b, c, 0, S32, 0x8771f681);
+	HH(c, d, a, b, 0, 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, 0, S34, 0xbebfbc70);
+	HH(a, b, c, d, 0, 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, 0, S31, 0xd9d4d039);
+	HH(d, a, b, c, 0, 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, 0, S41, 0x655b59c3);
+	II(d, a, b, c, x[3], S42, 0x8f0ccc92);
+	II(c, d, a, b, 0, S43, 0xffeff47d);
+	II(b, c, d, a, x[1], S44, 0x85845dd1);
+	II(a, b, c, d, 0, S41, 0x6fa87e4f);
+	II(d, a, b, c, 0, S42, 0xfe2ce6e0);
+	II(c, d, a, b, x[6], S43, 0xa3014314);
+	II(b, c, d, a, 0, S44, 0x4e0811a1);
+	II(a, b, c, d, x[4], S41, 0xf7537e82);
+	II(d, a, b, c, 0, S42, 0xbd3af235);
+	II(c, d, a, b, x[2], S43, 0x2ad7d2bb);
+	II(b, c, d, a, 0, S44, 0xeb86d391);
+
+	internal_ret[0] = a + (uint4) (0x67452301);
+	internal_ret[1] = b + (uint4) (0xefcdab89);
+	internal_ret[2] = c + (uint4) (0x98badcfe);
+	internal_ret[3] = d + (uint4) (0x10325476);
+}
+
+inline void clear_ctx(__private uint4 * x)
+{
+	int i;
+	uint4 zero = (uint4) (0, 0, 0, 0);
+	for (i = 0; i < 8; i++)
+		x[i] = zero;
+}
+inline void clean_ctx(__private uint *x){
+	int i;
+	for(i=0;i<8;i++);
+		x[i]=0;
+/*x++=0; nie dziala*/
+}
+
+
+__kernel void phpass
+    (__global const phpass_password * data, __global phpass_hash * data_out,
+    __global const char *setting) {
+	uint4 x[8];
+	clear_ctx(x);
+
+	uint idx = get_global_id(0);
+	uint4 a, b, c, d, x0, x1, x2, x3, x4, x5, x6, x7;
+
+	__global const uchar *password0 = data[idx * 4].v;
+	__global const uchar *password1 = data[idx * 4 + 1].v;
+	__global const uchar *password2 = data[idx * 4 + 2].v;
+	__global const uchar *password3 = data[idx * 4 + 3].v;
+	uchar4 length;		//password length
+	length.x = data[idx * 4 + 0].length;
+	length.y = data[idx * 4 + 1].length;
+	length.z = data[idx * 4 + 2].length;
+	length.w = data[idx * 4 + 3].length;
+	uint count, i = 0;
+	
+	uint sx[8];
+	__private uchar *buff = (uchar *) sx;
+ 	#define K(q)\
+		clean_ctx(sx);\
+		for (i = 0; i < 8; i++)\
+			buff[i] = setting[i];\
+		for (i = 8; i < 8 + length.s##q; i++)\
+			buff[i] = password##q[i - 8];\
+		for ( i = 0; i < 8; i++)\
+			x[i].s##q=sx[i];
+		K(0);
+		K(1);
+		K(2);
+		K(3);
+	#undef K
+
+	
+	uchar4 length2 = length + (uchar4) (8);
+	cuda_md5(length2, x, x);
+	#define K(q)\
+		clean_ctx(sx);\
+		for(i=0;i<length.s##q;i++)\
+			buff[i]=password##q[i];\
+		for(i=0;i<2;i++)\
+			x[i+4].s##q=sx[i];
+	K(0);
+	K(1);
+	K(2);
+	K(3);
+	#undef K
+
+
+	count = 1<< setting[SALT_SIZE + 3];
+
+
+	uchar4 length3 = length+(uchar4)(16);
+
+	uint4 x14 = (uint4)(length3.s0 << 3,length3.s1 << 3,length3.s2 << 3,length3.s3 << 3);
+
+	uint t;
+#define K(q)\
+	t=x[length3.s##q/4].s##q;\
+	t|= (((uint) 0x80) << ((length3.s##q & 0x3) << 3));\
+	x[length3.s##q/4].s##q=t;
+	K(0);
+	K(1);
+	K(2);
+	K(3);
+    
+#undef K
+//	x[length3.s0 / 4].s0 |= (((uint) 0x80) << ((length3.s0 & 0x3) << 3));
+//	x[length3.s1 / 4].s1 |= (((uint) 0x80) << ((length3.s1 & 0x3) << 3));
+//	x[length3.s2 / 4].s2 |= (((uint) 0x80) << ((length3.s2 & 0x3) << 3));
+//	x[length3.s3 / 4].s3 |= (((uint) 0x80) << ((length3.s3 & 0x3) << 3));
+
+
+	x0 = x[0];
+	x1 = x[1];
+	x2 = x[2];
+	x3 = x[3];
+	x4 = x[4];
+	x5 = x[5];
+	x6 = x[6];
+	x7 = x[7];
+	do {
+
+	   b = 0xefcdab89;
+	   c = 0x98badcfe;
+	   d = 0x10325476;
+
+	   //FF(a, b, c, d, x0, S11, 0xd76aa478);
+	  // FF(d, a, b, c, x1, S12, 0xe8c7b756);
+	  // FF(c, d, a, b, x2, S13, 0x242070db);
+	  // FF(b, c, d, a, x3, S14, 0xc1bdceee);
+	      a = AC1 + x0;
+		a = ROTATE_LEFT(a, S11);
+		a += b;		
+		d = (c ^ (a & MASK1)) + x1 + AC2pCd;
+		d = ROTATE_LEFT(d, S12);
+		d += a;		
+		c = F(d, a, b) + x2 + AC3pCc;
+		c = ROTATE_LEFT(c, S13);
+		c += d;		
+		b = F(c, d, a) + x3 + AC4pCb;
+		b = ROTATE_LEFT(b, S14);
+		b += c;
+ 	   FF(a, b, c, d, x4, S11, 0xf57c0faf);
+	   FF(d, a, b, c, x5, S12, 0x4787c62a);
+	   FF(c, d, a, b, x6, S13, 0xa8304613);
+	   FF(b, c, d, a, x7, S14, 0xfd469501);
+	  FF(a, b, c, d, 0, S11, 0x698098d8);
+	   FF(d, a, b, c, 0, S12, 0x8b44f7af);
+	   FF(c, d, a, b, 0, S13, 0xffff5bb1);
+	   FF(b, c, d, a, 0, S14, 0x895cd7be);
+	   FF(a, b, c, d, 0, S11, 0x6b901122);
+	   FF(d, a, b, c, 0, S12, 0xfd987193);
+	   FF(c, d, a, b, x14, S13, 0xa679438e);
+	   FF(b, c, d, a, 0, S14, 0x49b40821);
+
+	   GG(a, b, c, d, x1, S21, 0xf61e2562);
+	   GG(d, a, b, c, x6, S22, 0xc040b340);
+	   GG(c, d, a, b, 0, S23, 0x265e5a51);
+	   GG(b, c, d, a, x0, S24, 0xe9b6c7aa);
+	   GG(a, b, c, d, x5, S21, 0xd62f105d);
+	   GG(d, a, b, c, 0, S22, 0x2441453);
+	   GG(c, d, a, b, 0, S23, 0xd8a1e681);
+	   GG(b, c, d, a, x4, S24, 0xe7d3fbc8);
+	   GG(a, b, c, d, 0, S21, 0x21e1cde6);
+	   GG(d, a, b, c, x14, S22, 0xc33707d6);
+	   GG(c, d, a, b, x3, S23, 0xf4d50d87);
+	   GG(b, c, d, a, 0, S24, 0x455a14ed);
+	   GG(a, b, c, d, 0, S21, 0xa9e3e905);
+	   GG(d, a, b, c, x2, S22, 0xfcefa3f8);
+	   GG(c, d, a, b, x7, S23, 0x676f02d9);
+	   GG(b, c, d, a, 0, S24, 0x8d2a4c8a);
+
+	   HH(a, b, c, d, x5, S31, 0xfffa3942);
+	   HH(d, a, b, c, 0, S32, 0x8771f681);
+	   HH(c, d, a, b, 0, S33, 0x6d9d6122);
+	   HH(b, c, d, a, x14, S34, 0xfde5380c);
+	   HH(a, b, c, d, x1, S31, 0xa4beea44);
+	   HH(d, a, b, c, x4, S32, 0x4bdecfa9);
+	   HH(c, d, a, b, x7, S33, 0xf6bb4b60);
+	   HH(b, c, d, a, 0, S34, 0xbebfbc70);
+	   HH(a, b, c, d, 0, S31, 0x289b7ec6);
+	   HH(d, a, b, c, x0, S32, 0xeaa127fa);
+	   HH(c, d, a, b, x3, S33, 0xd4ef3085);
+	   HH(b, c, d, a, x6, S34, 0x4881d05);
+	   HH(a, b, c, d, 0, S31, 0xd9d4d039);
+	   HH(d, a, b, c, 0, S32, 0xe6db99e5);
+	   HH(c, d, a, b, 0, S33, 0x1fa27cf8);
+	   HH(b, c, d, a, x2, S34, 0xc4ac5665);
+
+	   II(a, b, c, d, x0, S41, 0xf4292244);
+	   II(d, a, b, c, x7, S42, 0x432aff97);
+	   II(c, d, a, b, x14, S43, 0xab9423a7);
+	   II(b, c, d, a, x5, S44, 0xfc93a039);
+	   II(a, b, c, d, 0, S41, 0x655b59c3);
+	   II(d, a, b, c, x3, S42, 0x8f0ccc92);
+	   II(c, d, a, b, 0, S43, 0xffeff47d);
+	   II(b, c, d, a, x1, S44, 0x85845dd1);
+	   II(a, b, c, d, 0, S41, 0x6fa87e4f);
+	   II(d, a, b, c, 0, S42, 0xfe2ce6e0);
+	   II(c, d, a, b, x6, S43, 0xa3014314);
+	   II(b, c, d, a, 0, S44, 0x4e0811a1);
+	   II(a, b, c, d, x4, S41, 0xf7537e82);
+	   II(d, a, b, c, 0, S42, 0xbd3af235);
+	   II(c, d, a, b, x2, S43, 0x2ad7d2bb);
+	   II(b, c, d, a, 0, S44, 0xeb86d391);
+
+	   x0 = a + (uint4)(0x67452301);
+	   x1 = b + (uint4)(0xefcdab89);
+	   x2 = c + (uint4)(0x98badcfe);
+	   x3 = d + (uint4)(0x10325476);
+
+	   } while (--count);
+
+
+
+
+	data_out[idx * 4 + 0].v[0] = x0.s0;
+	data_out[idx * 4 + 0].v[1] = x1.s0;
+	data_out[idx * 4 + 0].v[2] = x2.s0;
+	data_out[idx * 4 + 0].v[3] = x3.s0;
+
+	data_out[idx * 4 + 1].v[0] = x0.s1;
+	data_out[idx * 4 + 1].v[1] = x1.s1;
+	data_out[idx * 4 + 1].v[2] = x2.s1;
+	data_out[idx * 4 + 1].v[3] = x3.s1;
+
+	data_out[idx * 4 + 2].v[0] = x0.s2;
+	data_out[idx * 4 + 2].v[1] = x1.s2;
+	data_out[idx * 4 + 2].v[2] = x2.s2;
+	data_out[idx * 4 + 2].v[3] = x3.s2;
+
+	data_out[idx * 4 + 3].v[0] = x0.s3;
+	data_out[idx * 4 + 3].v[1] = x1.s3;
+	data_out[idx * 4 + 3].v[2] = x2.s3;
+	data_out[idx * 4 + 3].v[3] = x3.s3;
+
+}
diff -urpN john-1.7.9.orig//src/phpassopencl_fmt.c john-1.7.9-opencl_2//src/phpassopencl_fmt.c
--- john-1.7.9.orig//src/phpassopencl_fmt.c	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.9-opencl_2//src/phpassopencl_fmt.c	2011-12-18 04:16:32.062269044 +0000
@@ -0,0 +1,433 @@
+/*
+* 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"
+
+#ifdef __APPLE__
+#include <OpenCL/opencl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+#include <assert.h>
+#include "common_opencl.h"
+
+#define uint32_t		unsigned int
+#define uint16_t		unsigned short
+#define uint8_t			unsigned char
+
+#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 address(j,idx)		(((j)*KEYS_PER_CRYPT)+(idx))
+
+#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 SALT_SIZE		8
+
+#define THREADS 		256
+#define BLOCKS 			1
+//256*18
+#define KEYS_PER_CRYPT 		THREADS*BLOCKS*4
+#define MIN_KEYS_PER_CRYPT	KEYS_PER_CRYPT
+#define MAX_KEYS_PER_CRYPT	KEYS_PER_CRYPT
+#define FORMAT_LABEL		"phpassopencl"
+#define FORMAT_NAME		"PHPASSOPENCL"
+
+#define _PHPASS_DEBUG
+
+typedef struct {
+	unsigned char v[PLAINTEXT_LENGTH];
+	unsigned char length;
+} phpass_password;
+
+typedef struct {
+	uint32_t v[4];		///128bits for hash
+} phpass_hash;
+
+static phpass_password inbuffer[MAX_KEYS_PER_CRYPT];			/** plaintext ciphertexts **/
+static phpass_hash outbuffer[MAX_KEYS_PER_CRYPT];			/** calculated hashes **/
+static const char phpass_prefix[] = "$P$";
+
+static char currentsalt[SALT_SIZE + 1];
+
+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)
+{
+	opencl_init("phpass2.cl");
+}
+
+static void phpass_gpu(char *setting, int count_log2)
+{
+	/// Alocate memory
+	size_t insize = sizeof(phpass_password) * KEYS_PER_CRYPT;
+	size_t outsize = sizeof(phpass_hash) * KEYS_PER_CRYPT;
+	size_t settingsize = sizeof(uint8_t) * SALT_SIZE + 4;
+
+	cl_int cl_error;
+	cl_mem mem_in, mem_out, mem_setting;
+	mem_in =
+	    clCreateBuffer(context, CL_MEM_READ_ONLY, insize, NULL, &cl_error);
+	HANDLE_ERROR(cl_error);
+	mem_setting =
+	    clCreateBuffer(context, CL_MEM_READ_ONLY, settingsize, NULL,
+	    &cl_error);
+	HANDLE_ERROR(cl_error);
+	mem_out =
+	    clCreateBuffer(context, CL_MEM_WRITE_ONLY, outsize, NULL,
+	    &cl_error);
+	HANDLE_ERROR(cl_error);
+
+	/// Setup kernel parameters
+	cl_kernel kernel = clCreateKernel(prog, "phpass", &cl_error);
+	HANDLE_ERROR(cl_error);
+	clSetKernelArg(kernel, 0, sizeof(mem_in), &mem_in);
+	clSetKernelArg(kernel, 1, sizeof(mem_out), &mem_out);
+	clSetKernelArg(kernel, 2, sizeof(mem_setting), &mem_setting);
+
+
+	/// Copy data to gpu
+	HANDLE_ERROR(clEnqueueWriteBuffer(cq, mem_in, CL_FALSE, 0, insize,
+		inbuffer, 0, NULL, NULL));
+	HANDLE_ERROR(clEnqueueWriteBuffer(cq, mem_setting, CL_FALSE, 0,
+		settingsize, setting, 0, NULL, NULL));
+
+	/// Run kernel
+	size_t worksize = THREADS*BLOCKS;
+	size_t localworksize = THREADS;
+	HANDLE_ERROR(clEnqueueNDRangeKernel(cq, kernel, 1, NULL,
+		&worksize, &localworksize, 0, NULL, NULL));
+	HANDLE_ERROR(clFinish(cq));
+
+	/// Read the result back
+	HANDLE_ERROR(clEnqueueReadBuffer(cq, mem_out, CL_FALSE, 0, outsize,
+		outbuffer, 0, NULL, NULL));
+	/// Await completion of all the above
+	HANDLE_ERROR(clFinish(cq));
+
+	/// Cleanup
+	HANDLE_ERROR(clReleaseKernel(kernel));
+	HANDLE_ERROR(clReleaseMemObject(mem_in));
+	HANDLE_ERROR(clReleaseMemObject(mem_setting));
+	HANDLE_ERROR(clReleaseMemObject(mem_out));
+}
+
+static int valid(char *ciphertext)
+{
+	uint32_t i, j, count_log2, found;
+
+	if (strlen(ciphertext) != CIPHERTEXT_LENGTH)
+		return 0;
+	if (strncmp(ciphertext, phpass_prefix, 3) != 0)
+		return 0;
+
+	for (i = 3; i < CIPHERTEXT_LENGTH; i++) {
+		found = 0;
+		for (j = 0; j < 64; j++)
+			if (itoa64[j] == ARCH_INDEX(ciphertext[i])) {
+				found = 1;
+				break;
+			}
+		if (!found)
+			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 + 1];
+	memcpy(salt, &ciphertext[4], 8);
+	salt[8] = ciphertext[3];
+	return salt;
+}
+
+
+static void set_salt(void *salt)
+{
+	memcpy(currentsalt, salt, SALT_SIZE + 1);
+}
+
+static void set_key(char *key, int index)
+{
+	int 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];
+	memcpy(ret, inbuffer[index].v, inbuffer[index].length);
+	ret[inbuffer[index].length] = 0;
+	return ret;
+}
+
+static void crypt_all(int count)
+{
+	//setting format: salt+prefix+count_log2
+	char setting[SALT_SIZE + 3 + 1] = { 0 };
+	int count_log2 = 0;
+	count_log2 = atoi64[ARCH_INDEX(currentsalt[8])];
+	strcpy(setting, currentsalt);
+	strcpy(setting + SALT_SIZE, phpass_prefix);
+	setting[SALT_SIZE + 3] = count_log2;
+	phpass_gpu(setting, count_log2);
+}
+
+static int binary_hash_0(void *binary)
+{
+#ifdef _PHPASS_DEBUG
+	puts("binary_hash_0");
+	int i;
+	uint32_t *b = binary;
+	for (i = 0; i < 4; i++)
+		printf("%08x ", b[i]);
+	puts("");
+#endif
+	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 int get_hash_0(int index)
+{
+#ifdef _PHPASS_DEBUG
+	puts("get_hash");
+	int i;
+	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;
+	uint32_t b = ((uint32_t *) binary)[0];
+	//uint32_t *bb=binary;
+	//for(i=0;i<count;i++){
+	// printf("%x%x%x%x - %x%x%x%x",bb[0],bb[1],bb[2],bb[3],outbuffer[i].v[0],outbuffer[i].v[1],outbuffer[i].v[2],outbuffer[i].v[3]);   
+	//}
+	for (i = 0; i < count; i++)
+		if (b == outbuffer[i].v[0])
+			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[index].v[i])
+			return 0;
+	return 1;
+}
+
+static int cmp_exact(char *source, int count)
+{
+	return 1;
+}
+
+struct fmt_main fmt_PHPASSOPENCL = {
+	{
+		    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}
+};
