Openwall GNU/*/Linux - a small security-enhanced Linux distro for servers
[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Date: Fri, 9 Mar 2012 06:19:56 +0100
From: Lukas Odzioba <lukas.odzioba@...il.com>
To: john-dev@...ts.openwall.com
Subject: Re: CUDA & OpenCL status

It turned out that in magnum tree there were 2 versions od
phpass-opencl and cryptmd5-opencl files, old one and new one. Only old
one was working (new files were just lying there). I made a fix for
that.
magnum could you please merge it to your tree?

I think that kernels inside opencl dir now should not have opencl in
name, it was needed previously to have all opencl files in one place
on display list, now we should get rid of them.

I also added bitselect to phpass and cryptmd5 thanks Alexander.

Lukas

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}
-};

Powered by blists - more mailing lists

Your e-mail address:

Powered by Openwall GNU/*/Linux - Powered by OpenVZ