diff -urpN magnumripper-magnum-jumbo-3279fdc//src/Makefile magnumripper-magnum-jumbo-3279fdc-fixed//src/Makefile
--- magnumripper-magnum-jumbo-3279fdc//src/Makefile	2012-03-08 01:30:04.000000000 +0000
+++ magnumripper-magnum-jumbo-3279fdc-fixed//src/Makefile	2012-03-09 04:41:48.081304820 +0000
@@ -118,7 +118,7 @@ JOHN_OBJS = \
 
 OCL_OBJS = \
 	common-opencl.o opencl_mysqlsha1_fmt.o \
-	cryptmd5_opencl_fmt.o phpass_opencl_fmt.o opencl_rawsha1_fmt.o \
+	opencl_cryptmd5_fmt.o opencl_phpass_fmt.o opencl_rawsha1_fmt.o \
 	opencl_nt_fmt.o opencl_rawmd5_fmt.o  opencl_nsldaps_fmt.o
 
 CUDA_OBJS = \
diff -urpN magnumripper-magnum-jumbo-3279fdc//src/cryptmd5_opencl_fmt.c magnumripper-magnum-jumbo-3279fdc-fixed//src/cryptmd5_opencl_fmt.c
--- magnumripper-magnum-jumbo-3279fdc//src/cryptmd5_opencl_fmt.c	2012-03-08 01:30:04.000000000 +0000
+++ magnumripper-magnum-jumbo-3279fdc-fixed//src/cryptmd5_opencl_fmt.c	1970-01-01 00:00:00.000000000 +0000
@@ -1,498 +0,0 @@
-/*
-* 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"
-
-#include "common-opencl.h"
-#define uint32_t unsigned int
-#define uint8_t unsigned char
-
-#define KEYS_PER_CRYPT 		1024*9
-#define PLAINTEXT_LENGTH	15
-
-#define MIN(a,b) 		((a)<(b)?(a):(b))
-#define MAX(a,b) 		((a)>(b)?(a):(b))
-
-#define FORMAT_LABEL		"cryptmd5-opencl"
-#define FORMAT_NAME		"CRYPTMD5-OPENCL"
-#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$";
-//OpenCL variables:
-static cl_mem mem_in, mem_out, mem_salt;
-static size_t insize = sizeof(crypt_md5_password) * KEYS_PER_CRYPT;
-static size_t outsize = sizeof(crypt_md5_hash) * KEYS_PER_CRYPT;
-static size_t saltsize = sizeof(crypt_md5_salt);
-static size_t global_work_size = KEYS_PER_CRYPT;
-
-
-//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 release_all(void)
-{
-	HANDLE_CLERROR(clReleaseKernel(crypt_kernel), "Release kernel");
-	HANDLE_CLERROR(clReleaseMemObject(mem_in), "Release memin");
-	HANDLE_CLERROR(clReleaseMemObject(mem_salt), "Release memsalt");
-	HANDLE_CLERROR(clReleaseMemObject(mem_out), "Release memout");
-	HANDLE_CLERROR(clReleaseCommandQueue(queue[gpu_id]), "Release Queue");
-}
-
-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 find_best_workgroup()
-{
-	cl_event myEvent;
-	cl_ulong startTime, endTime, kernelExecTimeNs = CL_ULONG_MAX;
-	size_t my_work_group = 1;
-	cl_int ret_code;
-	int i;
-	size_t max_group_size;
-
-	clGetDeviceInfo(devices[gpu_id], CL_DEVICE_MAX_WORK_GROUP_SIZE,
-	    sizeof(max_group_size), &max_group_size, NULL);
-	cl_command_queue queue_prof =
-	    clCreateCommandQueue(context[gpu_id], devices[gpu_id],
-	    CL_QUEUE_PROFILING_ENABLE,
-	    &ret_code);
-	//printf("Max Group Work Size %d\n",(int)max_group_size);
-	local_work_size = 1;
-
-	/// Set keys
-	char *pass = "aaaaaaaa";
-	for (i = 0; i < KEYS_PER_CRYPT; i++) {
-		set_key(pass, i);
-	}
-	/// Copy data to GPU
-	HANDLE_CLERROR(clEnqueueWriteBuffer
-	    (queue_prof, mem_in, CL_FALSE, 0, insize, inbuffer, 0, NULL, NULL),
-	    "Copy memin");
-	HANDLE_CLERROR(clEnqueueWriteBuffer(queue_prof, mem_salt, CL_FALSE, 0,
-		saltsize, &host_salt, 0, NULL, NULL), "Copy memsalt");
-
-	/// Find minimum time
-	for (my_work_group = 1; (int) my_work_group <= (int) max_group_size;
-	    my_work_group *= 2) {
-
-		size_t localworksize = my_work_group;
-		HANDLE_CLERROR(clEnqueueNDRangeKernel
-		    (queue_prof, crypt_kernel, 1, NULL, &global_work_size,
-			&localworksize, 0, NULL, &myEvent), "Set ND range");
-
-
-		HANDLE_CLERROR(clFinish(queue_prof), "clFinish error");
-		clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_SUBMIT,
-		    sizeof(cl_ulong), &startTime, NULL);
-		clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_END,
-		    sizeof(cl_ulong), &endTime, NULL);
-
-		if ((endTime - startTime) < kernelExecTimeNs) {
-			kernelExecTimeNs = endTime - startTime;
-			local_work_size = my_work_group;
-		}
-		//printf("%d time=%lld\n",(int) my_work_group, endTime-startTime);
-	}
-	//printf("Optimal Group work Size = %d\n",(int)local_work_size);
-	clReleaseCommandQueue(queue_prof);
-}
-
-static void init(struct fmt_main *pFmt)
-{
-	opencl_init("$JOHN/cryptmd5_opencl_kernel.cl", gpu_id, platform_id);
-
-	///Alocate memory on the GPU
-	cl_int cl_error;
-	mem_in =
-	    clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, insize, NULL,
-	    &cl_error);
-	mem_salt =
-	    clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, saltsize, NULL,
-	    &cl_error);
-	mem_out =
-	    clCreateBuffer(context[gpu_id], CL_MEM_WRITE_ONLY, outsize, NULL,
-	    &cl_error);
-	///Assign kernel parameters 
-	crypt_kernel = clCreateKernel(program[gpu_id], KERNEL_NAME, &cl_error);
-	clSetKernelArg(crypt_kernel, 0, sizeof(mem_in), &mem_in);
-	clSetKernelArg(crypt_kernel, 1, sizeof(mem_out), &mem_out);
-	clSetKernelArg(crypt_kernel, 2, sizeof(mem_salt), &mem_salt);
-
-	find_best_workgroup();
-	//atexit(release_all);
-}
-
-
-static int valid(char *ciphertext, struct fmt_main *pFmt)
-{
-	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 int binary_hash_5(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0xffffff;
-}
-
-static int binary_hash_6(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0x7ffffff;
-}
-
-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 crypt_all(int count)
-{
-	///Copy data to GPU memory
-	HANDLE_CLERROR(clEnqueueWriteBuffer
-	    (queue[gpu_id], mem_in, CL_FALSE, 0, insize, inbuffer, 0, NULL,
-		NULL), "Copy memin");
-	HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_salt, CL_FALSE,
-		0, saltsize, &host_salt, 0, NULL, NULL), "Copy memsalt");
-
-	///Run kernel
-	size_t worksize = KEYS_PER_CRYPT;
-	size_t localworksize = local_work_size;
-	HANDLE_CLERROR(clEnqueueNDRangeKernel
-	    (queue[gpu_id], crypt_kernel, 1, NULL, &worksize, &localworksize,
-		0, NULL, NULL), "Set ND range");
-	HANDLE_CLERROR(clEnqueueReadBuffer(queue[gpu_id], mem_out, CL_FALSE, 0,
-		outsize, outbuffer, 0, NULL, NULL), "Copy data back");
-
-	///Await completion of all the above
-	HANDLE_CLERROR(clFinish(queue[gpu_id]), "clFinish error");
-}
-
-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 get_hash_5(int index)
-{
-	return ((ARCH_WORD_32 *) outbuffer[index].v)[0] & 0xffffff;
-}
-
-static int get_hash_6(int index)
-{
-	return ((ARCH_WORD_32 *) outbuffer[index].v)[0] & 0x7ffffff;
-}
-
-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_opencl_cryptMD5 = {
-	{
-		    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,
-		    fmt_default_prepare,
-		    valid,
-		    fmt_default_split,
-		    binary,
-		    salt,
-		    {
-				binary_hash_0,
-				binary_hash_1,
-				binary_hash_2,
-				binary_hash_3,
-				binary_hash_4,
-				binary_hash_5,
-			binary_hash_6},
-		    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,
-				get_hash_5,
-			get_hash_6},
-		    cmp_all,
-		    cmp_one,
-	    cmp_exact}
-};
diff -urpN magnumripper-magnum-jumbo-3279fdc//src/opencl/opencl_cryptmd5_kernel.cl magnumripper-magnum-jumbo-3279fdc-fixed//src/opencl/opencl_cryptmd5_kernel.cl
--- magnumripper-magnum-jumbo-3279fdc//src/opencl/opencl_cryptmd5_kernel.cl	2012-03-08 01:30:04.000000000 +0000
+++ magnumripper-magnum-jumbo-3279fdc-fixed//src/opencl/opencl_cryptmd5_kernel.cl	2012-03-09 05:08:16.263430126 +0000
@@ -8,8 +8,10 @@
 
 #define ROTATE_LEFT(x, s) rotate(x,s)
 
-#define F(x, y, z)	((z) ^ ((x) & ((y) ^ (z))))
-#define G(x, y, z)	((y) ^ ((z) & ((x) ^ (y))))
+//#define F(x, y, z)	((z) ^ ((x) & ((y) ^ (z))))
+//#define G(x, y, z)	((y) ^ ((z) & ((x) ^ (y))))
+#define F(x, y, z) bitselect((z), (y), (x))
+#define G(x, y, z) bitselect((y), (x), (z))
 
 #define H(x, y, z) (x^y^z)
 #define I(x, y, z) (y^(x|~z))
@@ -94,11 +96,8 @@ void ctx_update_private(__private md5_ct
 {
 	uint8_t *dest = &ctx->buffer[*ctx_buflen];
 	*ctx_buflen += len;
-	int i = len;
-	//while(len--)
-	 // *dest++=*string++;
-	for (i = 0; i < len; i++)
-		dest[i] = string[i];
+	while(len--)
+	  *dest++=*string++;
 }
 void ctx_insert_result(__private md5_ctx * ctx, __private uint8_t * string,uint8_t start)
 {
@@ -229,18 +228,6 @@ void md5_digest(__private md5_ctx * ctx,
 	result[3] = d + 0x10325476;
 }
 
-uint8_t next(int i, uint8_t saltlen,uint8_t passlen)
-{
-  uint8_t ret=0;
-  if ((i & 1) != 0){
-    if (i % 3 != 0)
-      ret+=saltlen;
-    if (i % 7 != 0)
-      ret+=passlen;
-  }
-  return ret;
-}
-
 __kernel void cryptmd5
     (__global const crypt_md5_password * inbuffer,
     __global uint32_t * outbuffer,
@@ -248,7 +235,6 @@ __kernel void cryptmd5
 	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;
@@ -282,77 +268,8 @@ __kernel void cryptmd5
 			ctx.buffer[ctx_buflen++] = pass[0];
 
 	md5_digest(&ctx, alt_result,&ctx_buflen);
-/*
-there are 8 cases:
-altpass
-altpasspass
-altsaltpass
-altsaltpasspass
-passalt
-passpassalt
-passsaltalt
-passsaltpassalt
-*/
-//prepare gtx
-	__private md5_ctx gtx[8];
-	__private alt_start[8];
-	uint8_t gtx_buflen[8];
-	for(i=0;i<4;i++) alt_start[i]=0;
-	for(i=0;i<8;i++) init_ctx(&gtx[i],&gtx_buflen[i]);
-	
-	{//altpass
-	  gtx_buflen[0]+=16;
-	  ctx_update_global(&gtx[0],(__global uint8_t *) pass, pass_len,&gtx_buflen[0]);
-	}
-	{//altpasspass
-	  gtx_buflen[1]+=16;
-	  ctx_update_global(&gtx[1],(__global uint8_t *) pass, pass_len,&gtx_buflen[1]);
-	  ctx_update_global(&gtx[1],(__global uint8_t *) pass, pass_len,&gtx_buflen[1]);
-	}
-	{//altsaltpass
-	  gtx_buflen[2]+=16;
-	  ctx_update_global(&gtx[2],(__global uint8_t *) salt, salt_len,&gtx_buflen[2]);
-	  ctx_update_global(&gtx[2],(__global uint8_t *) pass, pass_len,&gtx_buflen[2]);	
-	}
-	{//altsaltpasspass
-	  gtx_buflen[3]+=16;
-	  ctx_update_global(&gtx[3],(__global uint8_t *) salt, salt_len,&gtx_buflen[3]);
-	  ctx_update_global(&gtx[3],(__global uint8_t *) pass, pass_len,&gtx_buflen[3]);	
-	  ctx_update_global(&gtx[3],(__global uint8_t *) pass, pass_len,&gtx_buflen[3]);
-	}
-	{//passalt
-	  ctx_update_global(&gtx[4],(__global uint8_t *) pass, pass_len,&gtx_buflen[4]);
-	  gtx_buflen[4]+=16;
-	  alt_start[4]=pass_len;
-	}
-	{//passpassalt
-	  ctx_update_global(&gtx[5],(__global uint8_t *) pass, pass_len,&gtx_buflen[5]);
-	  ctx_update_global(&gtx[5],(__global uint8_t *) pass, pass_len,&gtx_buflen[5]);
-	  gtx_buflen[5]+=16;
-	  alt_start[5]=pass_len*2;
-	}
-	{//passsaltalt
-	  ctx_update_global(&gtx[6],(__global uint8_t *) pass, pass_len,&gtx_buflen[6]);	
-	  ctx_update_global(&gtx[6],(__global uint8_t *) salt, salt_len,&gtx_buflen[6]);
-	  gtx_buflen[6]+=16;
-	  alt_start[6]=pass_len+salt_len;
-	}
-	{//passsaltpassalt
-	  ctx_update_global(&gtx[7],(__global uint8_t *) pass, pass_len,&gtx_buflen[7]);	
-	  ctx_update_global(&gtx[7],(__global uint8_t *) salt, salt_len,&gtx_buflen[7]);
-	  ctx_update_global(&gtx[7],(__global uint8_t *) pass, pass_len,&gtx_buflen[7]);
-	  gtx_buflen[7]+=16;
-	  alt_start[7]=pass_len*2+salt_len;
-	}
-	uint8_t seq[]={0,7,3,5,3,7,1,6,3,5,3,7,1,7,2,5,3,7,1,7,3,4,3,7,1,7,3,5,2,7,1,7,3,5,3,6,1,7,3,5,3,7,7,3,5,3,7,1,6};
 
-	
-	for(i=0;i<1000;i++){
-	  int id=seq[i%42];//iteration id in gtx table
-	  ctx_insert_result(&gtx[id], (uint8_t*)alt_result, alt_start[id]);
-	  md5_digest(&gtx[id], alt_result,&gtx_buflen[id]);
-	}
-	/*for (i = 0; i < 1000; i++) {
+	for (i = 0; i < 1000; i++) {
 		init_ctx(&ctx,&ctx_buflen);
 		
 		if ((i & 1) != 0)
@@ -375,7 +292,7 @@ passsaltpassalt
 			ctx_update_global(&ctx, (__global uint8_t *) pass,
 			    pass_len,&ctx_buflen);
 		md5_digest(&ctx, alt_result,&ctx_buflen);
-	}*/
+	}
 	
 #define KEYS_PER_CRYPT 1024*9
 #define address(j,idx) 			(((j)*KEYS_PER_CRYPT)+(idx))
@@ -385,8 +302,4 @@ passsaltpassalt
 
 	K(0) K(1) K(2) K(3)
 	
-//	tresult[0] = alt_result[0];//ctx.A;
-//	tresult[1] = alt_result[1];//ctx.B;
-//	tresult[2] = alt_result[2];//ctx.C;
-//	tresult[3] = alt_result[3];//ctx.D;
 }
diff -urpN magnumripper-magnum-jumbo-3279fdc//src/opencl/opencl_phpass_kernel.cl magnumripper-magnum-jumbo-3279fdc-fixed//src/opencl/opencl_phpass_kernel.cl
--- magnumripper-magnum-jumbo-3279fdc//src/opencl/opencl_phpass_kernel.cl	2012-03-08 01:30:04.000000000 +0000
+++ magnumripper-magnum-jumbo-3279fdc-fixed//src/opencl/opencl_phpass_kernel.cl	2012-03-09 05:07:57.363429925 +0000
@@ -13,8 +13,11 @@ typedef struct {
 
 
 #define ROTATE_LEFT(x, s)	rotate(x,s)
-#define F(x, y, z)		((z) ^ ((x) & ((y) ^ (z))))
-#define G(x, y, z)		((y) ^ ((z) & ((x) ^ (y))))
+//#define F(x, y, z)		((z) ^ ((x) & ((y) ^ (z))))
+//#define G(x, y, z)		((y) ^ ((z) & ((x) ^ (y))))
+
+#define F(x, y, z) bitselect((z), (y), (x))
+#define G(x, y, z) bitselect((y), (x), (z))
 #define H(x, y, z)		((x) ^ (y) ^ (z))
 #define I(x, y, z)		((y) ^ ((x) | (~z)))
 
diff -urpN magnumripper-magnum-jumbo-3279fdc//src/opencl/phpass_opencl_kernel.cl magnumripper-magnum-jumbo-3279fdc-fixed//src/opencl/phpass_opencl_kernel.cl
--- magnumripper-magnum-jumbo-3279fdc//src/opencl/phpass_opencl_kernel.cl	2012-03-08 01:30:04.000000000 +0000
+++ magnumripper-magnum-jumbo-3279fdc-fixed//src/opencl/phpass_opencl_kernel.cl	1970-01-01 00:00:00.000000000 +0000
@@ -1,298 +0,0 @@
-
-#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) + (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 uint32_t unsigned int
-#define SALT_SIZE		8
-
-#define AC1				0xd76aa477
-#define AC2pCd				0xf8fa0bcc
-#define AC3pCc				0xbcdb4dd9
-#define AC4pCb				0xb18b7a77
-#define MASK1				0x77777777
-
-
-
-inline void cuda_md5(char len,__private uint32_t * internal_ret,__private 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;
-
-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 + 0x67452301;
-	internal_ret[1] = b + 0xefcdab89;
-	internal_ret[2] = c + 0x98badcfe;
-	internal_ret[3] = d + 0x10325476;
-}
-
-inline void clear_ctx(__private uint32_t * 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
-    )
-{
-	uint32_t x[8];
-	clear_ctx(x);
-	uint32_t a, b, c, d, x0, x1, x2, x3, x4, x5, x6, x7;
-
-	uint32_t idx = get_global_id(0);
-
-	__global const char *password = (__global const char*) data[idx].v;
-	int length, count, i;
-	__private unsigned char *buff = (unsigned char *) x;
-
-	length = data[idx].length;
-
-	for (i = 0; i < 8; i++)
-		buff[i] = setting[i];
-
-	for (i = 8; i < 8 + length; i++) {
-		buff[i] = password[i - 8];
-	}
-
-	cuda_md5(8 + length, x, x);
-	count = 1 << setting[SALT_SIZE+3];
-	for (i = 16; i < 16 + length; i++)
-		buff[i] = password[i - 16];
-
-
-	uint32_t len = 16 + length;
-	uint32_t x14 = len << 3;
-
-	x[len / 4] |= ((0x80) << ((len & 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);
-		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 + 0x67452301;
-		x1 = b + 0xefcdab89;
-		x2 = c + 0x98badcfe;
-		x3 = d + 0x10325476;
-
-	} while (--count);
-
-	data_out[idx].v[0] = x0;
-	data_out[idx].v[1] = x1;
-	data_out[idx].v[2] = x2;
-	data_out[idx].v[3] = x3;
-}
\ No newline at end of file
diff -urpN magnumripper-magnum-jumbo-3279fdc//src/opencl_cryptmd5_fmt.c magnumripper-magnum-jumbo-3279fdc-fixed//src/opencl_cryptmd5_fmt.c
--- magnumripper-magnum-jumbo-3279fdc//src/opencl_cryptmd5_fmt.c	2012-03-08 01:30:04.000000000 +0000
+++ magnumripper-magnum-jumbo-3279fdc-fixed//src/opencl_cryptmd5_fmt.c	2012-03-09 05:01:34.454179749 +0000
@@ -171,7 +171,7 @@ static void find_best_workgroup()
 	    clCreateCommandQueue(context[gpu_id], devices[gpu_id],
 	    CL_QUEUE_PROFILING_ENABLE,
 	    &ret_code);
-	//printf("Max Group Work Size %d\n",(int)max_group_size);
+	printf("Max Group Work Size %d\n",(int)max_group_size);
 	local_work_size = 1;
 
 	/// Set keys
@@ -208,13 +208,13 @@ static void find_best_workgroup()
 		}
 		//printf("%d time=%lld\n",(int) my_work_group, endTime-startTime);
 	}
-	//printf("Optimal Group work Size = %d\n",(int)local_work_size);
+	printf("Optimal Group work Size = %d\n",(int)local_work_size);
 	clReleaseCommandQueue(queue_prof);
 }
 
 static void init(struct fmt_main *pFmt)
 {
-	opencl_init("$JOHN/opencl_cryptmd5_kernel.cl", gpu_id);
+	opencl_init("$JOHN/opencl_cryptmd5_kernel.cl", gpu_id,platform_id);
 
 	///Alocate memory on the GPU
 
diff -urpN magnumripper-magnum-jumbo-3279fdc//src/opencl_nsldaps_fmt.c magnumripper-magnum-jumbo-3279fdc-fixed//src/opencl_nsldaps_fmt.c
--- magnumripper-magnum-jumbo-3279fdc//src/opencl_nsldaps_fmt.c	2012-03-08 01:30:04.000000000 +0000
+++ magnumripper-magnum-jumbo-3279fdc-fixed//src/opencl_nsldaps_fmt.c	2012-03-09 04:28:25.483272566 +0000
@@ -43,7 +43,7 @@
 #define NUM_BLOCKS			5
 
 #define PLAINTEXT_LENGTH		32
-#define SSHA_NUM_KEYS         		1024*2048*4
+#define SSHA_NUM_KEYS         		512*2048*4
 
 #define MIN_KEYS_PER_CRYPT              1024
 #define MAX_KEYS_PER_CRYPT		SSHA_NUM_KEYS
diff -urpN magnumripper-magnum-jumbo-3279fdc//src/opencl_phpass_fmt.c magnumripper-magnum-jumbo-3279fdc-fixed//src/opencl_phpass_fmt.c
--- magnumripper-magnum-jumbo-3279fdc//src/opencl_phpass_fmt.c	2012-03-08 01:30:04.000000000 +0000
+++ magnumripper-magnum-jumbo-3279fdc-fixed//src/opencl_phpass_fmt.c	2012-03-09 04:44:45.464429990 +0000
@@ -180,7 +180,7 @@ static void find_best_workgroup()
 static void init(struct fmt_main *pFmt)
 {
 	//atexit(release_all);
-	opencl_init("$JOHN/opencl_phpass_kernel.cl", gpu_id);
+	opencl_init("$JOHN/opencl_phpass_kernel.cl", gpu_id,platform_id);
 
 	/// Alocate memory
 	cl_int cl_error;
diff -urpN magnumripper-magnum-jumbo-3279fdc//src/phpass_opencl_fmt.c magnumripper-magnum-jumbo-3279fdc-fixed//src/phpass_opencl_fmt.c
--- magnumripper-magnum-jumbo-3279fdc//src/phpass_opencl_fmt.c	2012-03-08 01:30:04.000000000 +0000
+++ magnumripper-magnum-jumbo-3279fdc-fixed//src/phpass_opencl_fmt.c	1970-01-01 00:00:00.000000000 +0000
@@ -1,483 +0,0 @@
-/*
-* 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 <assert.h>
-#include "arch.h"
-#include "formats.h"
-#include "common.h"
-#include "misc.h"
-
-#include "common-opencl.h"
-
-#define uint32_t		unsigned int
-#define uint8_t			unsigned char
-
-#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 KEYS_PER_CRYPT		1024*9
-#define MIN_KEYS_PER_CRYPT	KEYS_PER_CRYPT
-#define MAX_KEYS_PER_CRYPT	KEYS_PER_CRYPT
-#define FORMAT_LABEL		"phpass-opencl"
-#define FORMAT_NAME		"PHPASS-OPENCL"
-
-//#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);
-
-// OpenCL variables:
-static cl_mem mem_in, mem_out, mem_setting;
-static size_t insize = sizeof(phpass_password) * KEYS_PER_CRYPT;
-static size_t outsize = sizeof(phpass_hash) * KEYS_PER_CRYPT;
-static size_t settingsize = sizeof(uint8_t) * SALT_SIZE + 4;
-static size_t global_work_size = KEYS_PER_CRYPT;
-
-
-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 release_all(void)
-{
-	HANDLE_CLERROR(clReleaseKernel(crypt_kernel), "Release Kernel");
-	HANDLE_CLERROR(clReleaseMemObject(mem_in), "Release mem in");
-	HANDLE_CLERROR(clReleaseMemObject(mem_setting), "Release mem setting");
-	HANDLE_CLERROR(clReleaseMemObject(mem_out), "Release mem out");
-	HANDLE_CLERROR(clReleaseCommandQueue(queue[gpu_id]), "Release Queue");
-}
-
-static void set_key(char *key, int index)
-{
-#ifdef _PHPASS_DEBUG
-	printf("set_key(%d) = %s\n", index, key);
-#endif
-	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 find_best_workgroup()
-{
-	cl_event myEvent;
-	cl_ulong startTime, endTime, kernelExecTimeNs = CL_ULONG_MAX;
-	size_t my_work_group = 1;
-	cl_int ret_code;
-	int i;
-	size_t max_group_size;
-	clGetDeviceInfo(devices[gpu_id], CL_DEVICE_MAX_WORK_GROUP_SIZE,
-	    sizeof(max_group_size), &max_group_size, NULL);
-	cl_command_queue queue_prof =
-	    clCreateCommandQueue(context[gpu_id], devices[gpu_id],
-	    CL_QUEUE_PROFILING_ENABLE,
-	    &ret_code);
-	HANDLE_CLERROR(ret_code, "Error while creating command queue");
-	local_work_size = 1;
-	/// Set keys
-	char *pass = "aaaaaaaa";
-	for (i = 0; i < KEYS_PER_CRYPT; i++) {
-		set_key(pass, i);
-	}
-	///Set salt
-	memcpy(currentsalt, "saltstri9", 9);
-	char setting[SALT_SIZE + 3 + 1] = { 0 };
-	strcpy(setting, currentsalt);
-	strcpy(setting + SALT_SIZE, phpass_prefix);
-	setting[SALT_SIZE + 3] = atoi64[ARCH_INDEX(currentsalt[8])];
-
-	///Copy data to GPU
-	HANDLE_CLERROR(clEnqueueWriteBuffer(queue_prof, mem_in, CL_FALSE, 0,
-		insize, inbuffer, 0, NULL, NULL), "Copy data to gpu");
-	HANDLE_CLERROR(clEnqueueWriteBuffer(queue_prof, mem_setting, CL_FALSE,
-		0, settingsize, setting, 0, NULL, NULL),
-	    "Copy setting to gpu");
-
-	///Find best local work size
-	for (my_work_group = 1; (int) my_work_group <= (int) max_group_size;
-	    my_work_group *= 2) {
-
-		HANDLE_CLERROR(clEnqueueNDRangeKernel(queue_prof, crypt_kernel,
-			1, NULL, &global_work_size, &my_work_group, 0, NULL,
-			&myEvent), "Run kernel");
-
-		HANDLE_CLERROR(clFinish(queue_prof), "clFinish error");
-		clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_SUBMIT,
-		    sizeof(cl_ulong), &startTime, NULL);
-		clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_END,
-		    sizeof(cl_ulong), &endTime, NULL);
-
-		if ((endTime - startTime) < kernelExecTimeNs) {
-			kernelExecTimeNs = endTime - startTime;
-			local_work_size = my_work_group;
-		}
-		//printf("%d time=%lld\n",(int) my_work_group, endTime-startTime);
-	}
-	printf("Optimal Group work Size = %d\n", (int) local_work_size);
-	clReleaseCommandQueue(queue_prof);
-}
-
-static void init(struct fmt_main *pFmt)
-{
-	//atexit(release_all);
-	opencl_init("$JOHN/phpass_opencl_kernel.cl", gpu_id, platform_id);
-
-	/// Alocate memory
-	cl_int cl_error;
-	mem_in =
-	    clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, insize, NULL,
-	    &cl_error);
-	HANDLE_CLERROR(cl_error, "Error alocating mem in");
-	mem_setting =
-	    clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, settingsize,
-	    NULL, &cl_error);
-	HANDLE_CLERROR(cl_error, "Error alocating mem setting");
-	mem_out =
-	    clCreateBuffer(context[gpu_id], CL_MEM_WRITE_ONLY, outsize, NULL,
-	    &cl_error);
-	HANDLE_CLERROR(cl_error, "Error alocating mem out");
-
-	/// Setup kernel parameters
-	crypt_kernel = clCreateKernel(program[gpu_id], "phpass", &cl_error);
-	HANDLE_CLERROR(cl_error, "Error creating kernel");
-	clSetKernelArg(crypt_kernel, 0, sizeof(mem_in), &mem_in);
-	clSetKernelArg(crypt_kernel, 1, sizeof(mem_out), &mem_out);
-	clSetKernelArg(crypt_kernel, 2, sizeof(mem_setting), &mem_setting);
-
-	find_best_workgroup();
-}
-
-static int valid(char *ciphertext, struct fmt_main *pFmt)
-{
-	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 crypt_all(int count)
-{
-#ifdef _PHPASS_DEBUG
-	printf("crypt_all(%d)\n", count);
-#endif
-	///Prepare setting format: salt+prefix+count_log2
-	char setting[SALT_SIZE + 3 + 1] = { 0 };
-	strcpy(setting, currentsalt);
-	strcpy(setting + SALT_SIZE, phpass_prefix);
-	setting[SALT_SIZE + 3] = atoi64[ARCH_INDEX(currentsalt[8])];
-	/// Copy data to gpu
-	HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_in, CL_FALSE, 0,
-		insize, inbuffer, 0, NULL, NULL), "Copy data to gpu");
-	HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_setting,
-		CL_FALSE, 0, settingsize, setting, 0, NULL, NULL),
-	    "Copy setting to gpu");
-
-	/// Run kernel
-	HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], crypt_kernel, 1,
-		NULL, &global_work_size, &local_work_size, 0, NULL, NULL),
-	    "Run kernel");
-	HANDLE_CLERROR(clFinish(queue[gpu_id]), "clFinish");
-
-	/// Read the result back
-	HANDLE_CLERROR(clEnqueueReadBuffer(queue[gpu_id], mem_out, CL_FALSE, 0,
-		outsize, outbuffer, 0, NULL, NULL), "Copy result back");
-
-	/// Await completion of all the above
-	HANDLE_CLERROR(clFinish(queue[gpu_id]), "clFinish");
-}
-
-static int binary_hash_0(void *binary)
-{
-#ifdef _PHPASS_DEBUG
-	printf("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 binary_hash_5(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0xffffff;
-}
-
-static int binary_hash_6(void *binary)
-{
-	return ((ARCH_WORD_32 *) binary)[0] & 0x7ffffff;
-}
-
-static int get_hash_0(int index)
-{
-#ifdef _PHPASS_DEBUG
-	printf("get_hash_0:   ");
-	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 get_hash_5(int index)
-{
-	return outbuffer[index].v[0] & 0xffffff;
-}
-
-static int get_hash_6(int index)
-{
-	return outbuffer[index].v[0] & 0x7ffffff;
-}
-
-static int cmp_all(void *binary, int count)
-{
-
-	uint32_t b = ((uint32_t *) binary)[0];
-	uint32_t i;
-	for (i = 0; i < count; i++) {
-		if (b == outbuffer[i].v[0]) {
-#ifdef _PHPASS_DEBUG
-			puts("cmp_all = 1");
-#endif
-			return 1;
-		}
-	}
-#ifdef _PHPASS_DEBUG
-	puts("cmp_all = 0");
-#endif	/* _PHPASS_DEBUG */
-	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]) {
-#ifdef _PHPASS_DEBUG
-			puts("cmp_one = 0");
-#endif
-			return 0;
-		}
-#ifdef _PHPASS_DEBUG
-	puts("cmp_one = 1");
-#endif
-	return 1;
-}
-
-static int cmp_exact(char *source, int count)
-{
-	return 1;
-}
-
-struct fmt_main fmt_opencl_phpass = {
-	{
-		    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,
-		    fmt_default_prepare,
-		    valid,
-		    fmt_default_split,
-		    binary,
-		    salt,
-		    {
-				binary_hash_0,
-				binary_hash_1,
-				binary_hash_2,
-				binary_hash_3,
-				binary_hash_4,
-				binary_hash_5,
-			binary_hash_6},
-		    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,
-				get_hash_5,
-			get_hash_6},
-		    cmp_all,
-		    cmp_one,
-	    cmp_exact}
-};