diff -urpN john-1.7.8.orig//src/Makefile opencl/john-1.7.8-cryptmd5opencl_0//src/Makefile
--- john-1.7.8.orig//src/Makefile	2011-05-04 18:52:31.000000000 +0000
+++ opencl/john-1.7.8-cryptmd5opencl_0//src/Makefile	2011-07-31 16:58:32.856804411 +0000
@@ -19,6 +19,7 @@ OMPFLAGS =
 #OMPFLAGS = -fopenmp
 # Sun Studio with OpenMP (set the OMP_NUM_THREADS env var at runtime)
 #OMPFLAGS = -xopenmp
+OCLHEADERPATH = /usr/local/cuda/include/
 CFLAGS = -c -Wall -O2 -fomit-frame-pointer $(OMPFLAGS)
 ASFLAGS = -c $(OMPFLAGS)
 LDFLAGS = -s $(OMPFLAGS)
@@ -33,6 +34,7 @@ JOHN_OBJS = \
 	AFS_fmt.o \
 	LM_fmt.o \
 	dummy.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 \
@@ -190,21 +192,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) -lcrypt -lOpenCL"
 
 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) -lcrypt -lOpenCL"
 
 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) -lcrypt -lOpenCL"
 
 linux-x86-avx:
 	$(LN) x86-sse.h arch.h
@@ -811,6 +813,9 @@ generic.h:
 		"$(BENCH_MD5_OBJS_DEPEND)" \
 		"$(BENCH_BF_OBJS_DEPEND)"
 
+cryptmd5opencl_fmt.o: cryptmd5opencl_fmt.c
+	$(CC) -I$(OCLHEADERPATH) $(CFLAGS) cryptmd5opencl_fmt.c
+
 bench: $(BENCH_OBJS)
 	$(LD) $(BENCH_OBJS) $(LDFLAGS) -o bench
 
diff -urpN john-1.7.8.orig//src/cryptmd5.cl opencl/john-1.7.8-cryptmd5opencl_0//src/cryptmd5.cl
--- john-1.7.8.orig//src/cryptmd5.cl	1970-01-01 00:00:00.000000000 +0000
+++ opencl/john-1.7.8-cryptmd5opencl_0//src/cryptmd5.cl	2011-07-31 19:28:04.305803963 +0000
@@ -0,0 +1,321 @@
+
+#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 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[ctx_buffsize];
+	uint32_t buflen;
+	uint32_t len;
+	uint32_t A, B, C, D;
+} md5_ctx;
+
+__constant char md5_salt_prefix[] = "$1$";
+__constant 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
+
+void ctx_update_global(__private md5_ctx * ctx, __global char *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 char *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 *dest=&ctx->buffer[ctx->buflen];
+	ctx->buflen+=3;
+	int i;
+	for(i=0;i<3;i++)
+	  dest[i]=md5_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 = & 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 char *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;
+	__global char *salt = hsalt->salt;
+
+	__private md5_ctx ctx, alt_ctx;
+	init_ctx(&ctx);
+	init_ctx(&alt_ctx);
+
+	ctx_update_global(&ctx, pass, pass_len);
+	ctx_update_prefix(&ctx);
+	ctx_update_global(&ctx, salt, salt_len);
+	
+	ctx_update_global(&alt_ctx, pass, pass_len);
+	ctx_update_global(&alt_ctx, salt, salt_len);
+	ctx_update_global(&alt_ctx, pass, pass_len);
+	md5_digest(&alt_ctx, alt_result);
+
+	for (i = pass_len; i > 16; i -= 16)
+		ctx_update_private(&ctx, alt_result, 16);
+	ctx_update_private(&ctx, 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, pass, pass_len);
+		else
+			ctx_update_private(&ctx, (const char *) alt_result, 16);
+
+		if (i % 3 != 0)
+			ctx_update_global(&ctx, salt, salt_len);
+
+		if (i % 7 != 0)
+			ctx_update_global(&ctx, pass, pass_len);
+
+		if ((i & 1) != 0)
+			ctx_update_private(&ctx, (const char *) alt_result, 16);
+		else
+			ctx_update_global(&ctx, 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.8.orig//src/cryptmd5.h opencl/john-1.7.8-cryptmd5opencl_0//src/cryptmd5.h
--- john-1.7.8.orig//src/cryptmd5.h	1970-01-01 00:00:00.000000000 +0000
+++ opencl/john-1.7.8-cryptmd5opencl_0//src/cryptmd5.h	2011-07-31 19:30:24.676804227 +0000
@@ -0,0 +1,47 @@
+/*
+* 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 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$";
+
+
+#endif
diff -urpN john-1.7.8.orig//src/cryptmd5opencl_fmt.c opencl/john-1.7.8-cryptmd5opencl_0//src/cryptmd5opencl_fmt.c
--- john-1.7.8.orig//src/cryptmd5opencl_fmt.c	1970-01-01 00:00:00.000000000 +0000
+++ opencl/john-1.7.8-cryptmd5opencl_0//src/cryptmd5opencl_fmt.c	2011-07-31 18:37:25.950804149 +0000
@@ -0,0 +1,461 @@
+/*
+* 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"
+#include <CL/cl.h>
+
+
+#define FORMAT_LABEL		"cryptmd5opencl"
+#define FORMAT_NAME		"CRYPTMD5OPENCL"
+
+#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
+
+static 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 **/
+
+#define SOURCE_BUFSIZE 1024*64
+
+static char phpass_cl[SOURCE_BUFSIZE];
+static cl_context context;
+static cl_command_queue cq;
+static cl_program prog;
+static size_t src_size = 0;
+
+
+//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 md5_crypt_gpu(crypt_md5_password * inbuffer,
+    crypt_md5_hash * outbuffer, crypt_md5_salt * host_salt){
+
+	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);
+ 
+	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);
+	
+	cl_kernel phpass_kernel = clCreateKernel(prog, "cryptmd5", &cl_error);
+	clSetKernelArg(phpass_kernel, 0, sizeof(mem_in), &mem_in);
+	clSetKernelArg(phpass_kernel, 1, sizeof(mem_out), &mem_out);
+	clSetKernelArg(phpass_kernel, 2, sizeof(mem_salt), &mem_salt);
+	//copy data to gpu
+	cl_error =
+	    clEnqueueWriteBuffer(cq, mem_in, CL_FALSE, 0, insize, inbuffer, 0, NULL,
+	    NULL);
+	cl_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;
+	cl_error =
+	    clEnqueueNDRangeKernel(cq, phpass_kernel, 1, NULL, &worksize,
+	    &localworksize, 0, NULL, NULL);
+	cl_error =
+	    clEnqueueReadBuffer(cq, mem_out, CL_FALSE, 0, outsize, outbuffer, 0, NULL,
+	    NULL);
+	// Await completion of all the above
+	cl_error = clFinish(cq);    
+}
+
+static void readSource(char *src, const char *path)
+{
+	FILE *f = fopen(path, "r");
+	if (f != NULL) {
+		src_size = fread(src, 1, SOURCE_BUFSIZE, f);
+		fclose(f);
+	}
+}
+
+static void init(void)
+{
+	readSource(phpass_cl, "cryptmd5.cl");
+	cl_int cl_error;
+	cl_platform_id platform;
+	cl_device_id device;
+	cl_uint platforms, devices;
+	cl_error = clGetPlatformIDs(1, &platform, &platforms);
+	cl_error =
+	    clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, &devices);
+	cl_context_properties properties[] = {
+		CL_CONTEXT_PLATFORM, (cl_context_properties) platform,
+		0
+	};
+	//
+	char log[1024*64];
+	cl_error=clGetDeviceInfo(device,CL_DEVICE_NAME,sizeof(log),log,NULL);
+	printf("\nOpenCL Device: %s\n",log);
+	
+	//end
+	context =
+	    clCreateContext(properties, 1, &device, NULL, NULL, &cl_error);
+	cq = clCreateCommandQueue(context, device, 0, &cl_error);
+	const char *srcptr[] = { phpass_cl };
+	prog = clCreateProgramWithSource(context,
+	    1, srcptr, &src_size, &cl_error);
+	if (cl_error != CL_SUCCESS) {
+		printf("Error while creating program [%d]\n", cl_error);
+		exit(1);
+	}
+	cl_error = clBuildProgram(prog, 0, NULL, "", NULL, NULL);
+	if (cl_error != CL_SUCCESS) {
+		printf("Error while building program [%d]\n", cl_error);
+		cl_build_status build_status;
+		cl_error =
+		    clGetProgramBuildInfo(prog, device,
+		    CL_PROGRAM_BUILD_STATUS, sizeof(cl_build_status),
+		    &build_status, NULL);
+
+		char *build_log;
+		size_t ret_val_size;
+		cl_error =
+		    clGetProgramBuildInfo(prog, device, CL_PROGRAM_BUILD_LOG,
+		    0, NULL, &ret_val_size);
+
+		build_log = malloc(ret_val_size + 1);
+		cl_error =
+		    clGetProgramBuildInfo(prog, device, CL_PROGRAM_BUILD_LOG,
+		    ret_val_size, build_log, NULL);
+		build_log[ret_val_size] = '\0';
+		printf("BUILD LOG: \n %s", build_log);
+		free(build_log);
+		//exit(1);
+	}
+}
+
+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)
+{
+/*	puts("binary_hash");
+	int i;
+	uint32_t *b=binary;
+	for(i=0;i<4;i++)
+	  printf("%08x ",b[i]);
+	puts("");
+*/	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)
+{
+//	printf("set_key(%s,%d)\n",key,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)
+{
+/*	int i,j;
+	puts("inbuffer");
+	for(i=0;i<count;i++){
+	 printf("%d =",inbuffer[i].length);
+	 for(j=0;j<inbuffer[i].length;j++)
+	   printf("%c",inbuffer[i].v[j]);
+	 puts("");
+	  
+	}
+*/	md5_crypt_gpu(inbuffer, outbuffer, &host_salt);
+	
+}
+
+static int get_hash_0(int index)
+{
+ /* puts("get_hash_0");
+  int i;
+  for(i=0;i<4;i++)
+    printf("%08x ",outbuffer[index].v[i]);
+   puts("");
+*/	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.8.orig//src/john.c opencl/john-1.7.8-cryptmd5opencl_0//src/john.c
--- john-1.7.8.orig//src/john.c	2011-02-27 12:31:36.000000000 +0000
+++ opencl/john-1.7.8-cryptmd5opencl_0//src/john.c	2011-07-31 16:50:51.307857975 +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_CRYPTMD5OPENCL;
 #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_CRYPTMD5OPENCL);
 #ifdef HAVE_CRYPT
 	john_register_one(&fmt_crypt);
 #endif
diff -urpN john-1.7.8.orig//src/options.c opencl/john-1.7.8-cryptmd5opencl_0//src/options.c
--- john-1.7.8.orig//src/options.c	2011-06-22 13:01:40.000000000 +0000
+++ opencl/john-1.7.8-cryptmd5opencl_0//src/options.c	2011-07-31 16:50:09.581890957 +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/CRYPTMD5OPENCL" MAYBE_CRYPT "\n" \
 "--save-memory=LEVEL        enable memory saving, at LEVEL 1..3\n"
 
 void opt_init(char *name, int argc, char **argv)
