diff -urpN john-1.7.6_jumbo7/run/nt_opencl_kernel.c john-1.7.6_ntopencl/run/nt_opencl_kernel.c
--- john-1.7.6_jumbo7/run/nt_opencl_kernel.c	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.6_ntopencl/run/nt_opencl_kernel.c	2010-09-05 00:47:39.000000000 +0000
@@ -0,0 +1,111 @@
+//Init values
+#define INIT_A 0x67452301
+#define INIT_B 0xefcdab89
+#define INIT_C 0x98badcfe
+#define INIT_D 0x10325476
+
+#define SQRT_2 0x5a827999
+#define SQRT_3 0x6ed9eba1
+
+#define PLAINTEXT_LENGTH 27
+
+__kernel void nt_crypt(const __global char *keys , __global uint *bbbs,__global uint *nt_buffer1x,__global uint *output1x)
+{
+	uint i = get_global_id(0);
+	__global char *key=keys+(i*32);
+	__global uint *nt_buffer=nt_buffer1x+i*16;
+	
+	//set ket
+	uint nt_index = 0;
+	uint md4_size = 0;
+	
+	for(;key[md4_size] && md4_size<PLAINTEXT_LENGTH;nt_index++,md4_size++)
+	{
+		ushort temp = (ushort)(key[++md4_size]);
+		nt_buffer[nt_index] = upsample( (ushort)(temp ? temp : 0x80), (ushort)(key[md4_size-1]));
+
+		if(!temp)
+			goto key_cleaning;
+
+	}
+	
+	nt_buffer[nt_index] = 0x80;
+	
+key_cleaning:
+	nt_index++;
+	for(;nt_index < 14; nt_index++)
+		nt_buffer[nt_index] = 0;
+	
+	nt_buffer[14] = md4_size << 4;
+	//end set key
+	
+	uint a;
+	uint b;
+	uint c;
+	uint d;
+	
+	/* Round 1 */
+	a = 		0xFFFFFFFF					 + nt_buffer[0]; a=rotate(a, (uint)3);
+	d = INIT_D+(INIT_C ^ (a & 0x77777777))   + nt_buffer[1]; d=rotate(d, (uint)7);
+	c = INIT_C+(INIT_B ^ (d & (a ^ INIT_B))) + nt_buffer[2]; c=rotate(c, (uint)11);
+	b = INIT_B + (a ^ (c & (d ^ a)))		 + nt_buffer[3]; b=rotate(b, (uint)19);
+	
+	a += (d ^ (b & (c ^ d)))  +  nt_buffer[4] ; a = rotate(a , (uint)3 );
+	d += (c ^ (a & (b ^ c)))  +  nt_buffer[5] ; d = rotate(d , (uint)7 );
+	c += (b ^ (d & (a ^ b)))  +  nt_buffer[6] ; c = rotate(c , (uint)11);
+	b += (a ^ (c & (d ^ a)))  +  nt_buffer[7] ; b = rotate(b , (uint)19);
+	
+	a += (d ^ (b & (c ^ d)))  +  nt_buffer[8] ; a = rotate(a , (uint)3 );
+	d += (c ^ (a & (b ^ c)))  +  nt_buffer[9] ; d = rotate(d , (uint)7 );
+	c += (b ^ (d & (a ^ b)))  +  nt_buffer[10]; c = rotate(c , (uint)11);
+	b += (a ^ (c & (d ^ a)))  +  nt_buffer[11]; b = rotate(b , (uint)19);
+	
+	a += (d ^ (b & (c ^ d)))  +  nt_buffer[12]; a = rotate(a , (uint)3 );
+	d += (c ^ (a & (b ^ c)))  +  nt_buffer[13]; d = rotate(d , (uint)7 );
+	c += (b ^ (d & (a ^ b)))  +  nt_buffer[14]; c = rotate(c , (uint)11);
+	b += (a ^ (c & (d ^ a)))                  ; b = rotate(b , (uint)19);
+	
+	/* Round 2 */
+	a += ((b & (c | d)) | (c & d)) + nt_buffer[0] + SQRT_2; a = rotate(a , (uint)3 );
+	d += ((a & (b | c)) | (b & c)) + nt_buffer[4] + SQRT_2; d = rotate(d , (uint)5 );
+	c += ((d & (a | b)) | (a & b)) + nt_buffer[8] + SQRT_2; c = rotate(c , (uint)9 );
+	b += ((c & (d | a)) | (d & a)) + nt_buffer[12]+ SQRT_2; b = rotate(b , (uint)13);
+	
+	a += ((b & (c | d)) | (c & d)) + nt_buffer[1] + SQRT_2; a = rotate(a , (uint)3 );
+	d += ((a & (b | c)) | (b & c)) + nt_buffer[5] + SQRT_2; d = rotate(d , (uint)5 );
+	c += ((d & (a | b)) | (a & b)) + nt_buffer[9] + SQRT_2; c = rotate(c , (uint)9 );
+	b += ((c & (d | a)) | (d & a)) + nt_buffer[13]+ SQRT_2; b = rotate(b , (uint)13);
+	
+	a += ((b & (c | d)) | (c & d)) + nt_buffer[2] + SQRT_2; a = rotate(a , (uint)3 );
+	d += ((a & (b | c)) | (b & c)) + nt_buffer[6] + SQRT_2; d = rotate(d , (uint)5 );
+	c += ((d & (a | b)) | (a & b)) + nt_buffer[10]+ SQRT_2; c = rotate(c , (uint)9 );
+	b += ((c & (d | a)) | (d & a)) + nt_buffer[14]+ SQRT_2; b = rotate(b , (uint)13);
+	
+	a += ((b & (c | d)) | (c & d)) + nt_buffer[3] + SQRT_2; a = rotate(a , (uint)3 );
+	d += ((a & (b | c)) | (b & c)) + nt_buffer[7] + SQRT_2; d = rotate(d , (uint)5 );
+	c += ((d & (a | b)) | (a & b)) + nt_buffer[11]+ SQRT_2; c = rotate(c , (uint)9 );
+	b += ((c & (d | a)) | (d & a))                + SQRT_2; b = rotate(b , (uint)13);
+	
+	/* Round 3 */
+	a += (d ^ c ^ b) + nt_buffer[0]  + SQRT_3; a = rotate(a , (uint)3 );
+	d += (c ^ b ^ a) + nt_buffer[8]  + SQRT_3; d = rotate(d , (uint)9 );
+	c += (b ^ a ^ d) + nt_buffer[4]  + SQRT_3; c = rotate(c , (uint)11);
+	b += (a ^ d ^ c) + nt_buffer[12] + SQRT_3; b = rotate(b , (uint)15);
+	
+	a += (d ^ c ^ b) + nt_buffer[2]  + SQRT_3; a = rotate(a , (uint)3 );
+	d += (c ^ b ^ a) + nt_buffer[10] + SQRT_3; d = rotate(d , (uint)9 );
+	c += (b ^ a ^ d) + nt_buffer[6]  + SQRT_3; c = rotate(c , (uint)11);
+	b += (a ^ d ^ c) + nt_buffer[14] + SQRT_3; b = rotate(b , (uint)15);
+	
+	a += (d ^ c ^ b) + nt_buffer[1]  + SQRT_3; a = rotate(a , (uint)3 );
+	d += (c ^ b ^ a) + nt_buffer[9]  + SQRT_3; d = rotate(d , (uint)9 );
+	c += (b ^ a ^ d) + nt_buffer[5]  + SQRT_3; c = rotate(c , (uint)11);
+	b += (a ^ d ^ c) + nt_buffer[13];
+	
+	output1x[4*i+0] = a;
+	output1x[4*i+1] = b;
+	output1x[4*i+2] = c;
+	output1x[4*i+3] = d;
+	
+	bbbs[i] = b;
+}
diff -urpN john-1.7.6_jumbo7/src/Makefile john-1.7.6_ntopencl/src/Makefile
--- john-1.7.6_jumbo7/src/Makefile	2010-08-22 16:47:16.000000000 +0000
+++ john-1.7.6_ntopencl/src/Makefile	2010-09-05 00:17:51.000000000 +0000
@@ -182,6 +182,8 @@ default:
 	@echo "sco-x86-any-cc           SCO, x86, cc"
 	@echo "tru64-alpha              Tru64 (Digital UNIX, OSF/1), Alpha"
 	@echo "aix-ppc32                AIX, PowerPC 32-bit"
+	@echo "macosx-x86-64-opencl     Mac OS X 10.6+, Xcode 3.0+, x86-64 with SSE2 and OpenCL (experimental)"
+	@echo "macosx-x86-sse2-opencl   Mac OS X 10.6+, x86 with SSE2 and OpenCL (experimental)"
 	@echo "macosx-x86-64            Mac OS X 10.5+, Xcode 3.0+, x86-64 with SSE2 (best)"
 	@echo "macosx-x86-sse2          Mac OS X, x86 with SSE2"
 	@echo "macosx-ppc32-altivec     Mac OS X, PowerPC w/AltiVec (best)"
@@ -603,6 +605,22 @@ aix-ppc32:
 		OPT_NORMAL="-O2" \
 		OPT_INLINE="-O3 -Q=99 -w"
 
+macosx-x86-64-opencl:
+	$(LN) x86-64.h arch.h
+	$(MAKE) $(PROJ) \
+		JOHN_OBJS="$(JOHN_OBJS) x86-64.o" \
+		ASFLAGS="$(ASFLAGS) -m64 -DUNDERSCORES -DBSD -DALIGN_LOG" \
+		CFLAGS="$(CFLAGS) -m64 -DNT_OPENCL" \
+		LDFLAGS="$(LDFLAGS) -m64 -framework OpenCL"
+		
+macosx-x86-sse2-opencl:
+	$(LN) x86-sse.h arch.h
+	$(MAKE) $(PROJ) \
+		JOHN_OBJS="$(JOHN_OBJS) x86.o x86-sse.o sha1-mmx.o md5-mmx.o" \
+		ASFLAGS="$(ASFLAGS) -m32 -DUNDERSCORES -DBSD -DALIGN_LOG" \
+		CFLAGS="$(CFLAGS) -m32 -DNT_OPENCL" \
+		LDFLAGS="$(LDFLAGS) -m32 -framework OpenCL"
+		
 macosx-x86-64:
 	$(LN) x86-64.h arch.h
 	$(MAKE) $(PROJ) \
diff -urpN john-1.7.6_jumbo7/src/NT_fmt.c john-1.7.6_ntopencl/src/NT_fmt.c
--- john-1.7.6_jumbo7/src/NT_fmt.c	2010-08-22 16:22:31.000000000 +0000
+++ john-1.7.6_ntopencl/src/NT_fmt.c	2010-09-05 18:05:45.000000000 +0000
@@ -1,10 +1,10 @@
-/* NTLM patch for john (performance improvement)
+/* NTLM patch for john (performance improvement and OpenCL 1.0 conformant)
  *
- * Written by Alain Espinosa <alainesp at gmail.com> in 2007.  No copyright
+ * Written by Alain Espinosa <alainesp at gmail.com> in 2010.  No copyright
  * is claimed, and the software is hereby placed in the public domain.
  * In case this attempt to disclaim copyright and place the software in the
  * public domain is deemed null and void, then the software is
- * Copyright (c) 2007 Alain Espinosa and it is hereby released to the
+ * Copyright (c) 2010 Alain Espinosa and it is hereby released to the
  * general public under the following terms:
  *
  * Redistribution and use in source and binary forms, with or without
@@ -21,6 +21,11 @@
 #include "memory.h"
 #include "common.h"
 #include "formats.h"
+#include "path.h"
+#ifdef NT_OPENCL
+	#include <OpenCL/cl.h>
+#endif
+
 
 //Init values
 #define INIT_A 0x67452301
@@ -52,7 +57,42 @@ static struct fmt_tests tests[] = {
 #define BINARY_SIZE			16
 #define SALT_SIZE			0
 
-#if defined (NT_X86_64)
+#ifdef NT_OPENCL
+	#define NT_NUM_KEYS			1024*16
+
+	//Putting here for succeful compilation (Needed by assembly functions).
+	//Maybe useful in the future perform CPU and GPU cracking side by side
+	unsigned int *nt_buffer8x;
+	unsigned int *output8x;
+	unsigned int *nt_buffer4x;
+	unsigned int *output4x;
+	unsigned int *nt_buffer1x;
+	unsigned int *output1x;
+
+	static cl_uint bbbs[NT_NUM_KEYS];
+	static char saved_plain[32*NT_NUM_KEYS];
+
+	cl_command_queue queue;
+	cl_kernel nt_crypt_kernel;
+	cl_mem buffer_nt;
+	cl_mem buffer_out;
+	cl_mem buffer_keys;
+	cl_mem buffer_bbbs;
+	#define ALGORITHM_NAME		"OpenCL 1.0"
+	#define NT_CRYPT_FUN		nt_crypt_all_opencl
+	//Putting in this function error logs can degrade performance
+	static void nt_crypt_all_opencl(int count)
+	{
+		size_t global_work_size = NT_NUM_KEYS;
+		
+		//If you are using CPU device comment the following line for a small speeud
+		clEnqueueWriteBuffer(queue, buffer_keys, CL_TRUE, 0, sizeof(saved_plain), (void*)saved_plain, 0, NULL, NULL);
+		clEnqueueNDRangeKernel( queue, nt_crypt_kernel, 1, NULL, &global_work_size, NULL, 0, NULL, NULL);
+		clFinish( queue );
+		//If you are using CPU device comment the following line for a small speeud
+		clEnqueueReadBuffer(queue, buffer_bbbs, CL_TRUE, 0, sizeof(bbbs), (void*)bbbs, 0, NULL, NULL);
+	}
+#elif defined (NT_X86_64)
 	#define NT_NUM_KEYS	32
 	
 	unsigned int nt_buffer8x[16*NT_NUM_KEYS] __attribute__ ((aligned(16)));
@@ -158,24 +198,88 @@ static struct fmt_tests tests[] = {
 	}	
 #endif
 
-static unsigned int last_i[NT_NUM_KEYS];
-static char saved_plain[32*NT_NUM_KEYS];
-
 #define MIN_KEYS_PER_CRYPT		NT_NUM_KEYS
 #define MAX_KEYS_PER_CRYPT		NT_NUM_KEYS
 
-static void fmt_NT_init(void)
-{
-	memset(last_i,0,4*NT_NUM_KEYS);
-#if defined(NT_X86_64)
-	memset(nt_buffer8x,0,16*4*NT_NUM_KEYS);
-#elif defined(NT_SSE2)
-	memset(nt_buffer4x,0,64*4*NT_NUM_KEYS1);
-	memset(nt_buffer1x,0,16*4*NT_NUM_KEYS1);
+#ifdef NT_OPENCL
+	static void if_error_log(cl_int ret_code, char* message)
+	{
+		if (ret_code != CL_SUCCESS)
+			printf("\nOpenCL: %s\n",message);
+	}
+	static void fmt_NT_init(void)
+	{
+		FILE* file = NULL;
+		char log[1024*64];      // Build log
+		char *source;	        // Device Program Source
+		cl_int ret_code;
+		cl_platform_id platform;
+		cl_device_id device;
+		cl_context context;
+		cl_program program;
+		
+		// Load the source
+		source = (char*)mem_alloc(1024*16); // Assume max 16KB of source
+		file = fopen(path_expand("$JOHN/nt_opencl_kernel.c"),"r");                                                     if_error_log((file != NULL ? CL_SUCCESS : -1),"Source kernel not found");
+		size_t source_size = fread(source,sizeof(char),1024*16,file);
+		source[source_size] = 0;            //NUll terminate it
+		fclose(file);
+		
+		// 1. Get a platform.
+		ret_code = clGetPlatformIDs( 1, &platform, NULL );                                                             if_error_log (ret_code,"No OpenCL platform exist");
+		
+		// 2. Find an OpenCL device.
+		// Use CL_DEVICE_TYPE_CPU for a CPU device
+		// Use CL_DEVICE_TYPE_GPU for a Graphic Card
+		// Use CL_DEVICE_TYPE_ACCELERATOR if you have, for example, an IBM Cell Blade accelerator
+		ret_code = clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);                                    if_error_log (ret_code,"No OpenCL device of that type exist");
+		
+		// 3. Create a context and command queue on that device.
+		context = clCreateContext( NULL, 1, &device, NULL, NULL, &ret_code);                                           if_error_log (ret_code,"Error creating context");
+		queue = clCreateCommandQueue( context, device, 0, &ret_code );                                                 if_error_log (ret_code,"Error creating command queue");
+		
+		// 4. Perform runtime source compilation, and obtain kernel entry point.
+		program = clCreateProgramWithSource( context, 1, &source, NULL, &ret_code );                                   if_error_log (ret_code,"Error creating program");
+		ret_code = clBuildProgram( program, 1, &device, NULL, NULL, NULL );                                            if_error_log (ret_code,"Error building program. Log:");
+		clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, sizeof(log), (void*)log, NULL);                   if_error_log (ret_code,log);
+ 
+		// 5. Create kernel to execute
+		nt_crypt_kernel = clCreateKernel( program, "nt_crypt", &ret_code );                                            if_error_log (ret_code,"Error creating kernel");
+		
+		// 6. Create and set arguments
+		buffer_nt   = clCreateBuffer( context, CL_MEM_READ_WRITE, 16*sizeof(cl_uint)*NT_NUM_KEYS, NULL , &ret_code );  if_error_log (ret_code,"Error creating nt_crypt buffer");
+		buffer_out  = clCreateBuffer( context, CL_MEM_READ_WRITE, 4*sizeof(cl_uint)*NT_NUM_KEYS , NULL , &ret_code );  if_error_log (ret_code,"Error creating output buffer");
+		// If you are using a CPU device uncomment the following 2 lines for a small speedup
+		//buffer_keys = clCreateBuffer( context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, sizeof(saved_plain) , saved_plain , &ret_code );if_error_log (ret_code,"Error creating keys buffer");
+		//buffer_bbbs = clCreateBuffer( context, CL_MEM_WRITE_ONLY| CL_MEM_USE_HOST_PTR, sizeof(bbbs)        , bbbs        , &ret_code );if_error_log (ret_code,"Error creating bbbs buffer");
+		// For a general device
+		buffer_keys = clCreateBuffer( context, CL_MEM_READ_ONLY , sizeof(saved_plain)           , NULL , &ret_code );  if_error_log (ret_code,"Error creating keys buffer");
+		buffer_bbbs = clCreateBuffer( context, CL_MEM_WRITE_ONLY, sizeof(bbbs)                  , NULL , &ret_code );  if_error_log (ret_code,"Error creating bbbs buffer");
+		
+		ret_code = clSetKernelArg(nt_crypt_kernel, 0, sizeof(buffer_keys), (void*) &buffer_keys);                      if_error_log (ret_code,"Error setting argument 0");
+		ret_code = clSetKernelArg(nt_crypt_kernel, 1, sizeof(buffer_bbbs), (void*) &buffer_bbbs);                      if_error_log (ret_code,"Error setting argument 1");
+		ret_code = clSetKernelArg(nt_crypt_kernel, 2, sizeof(buffer_nt)  , (void*) &buffer_nt);                        if_error_log (ret_code,"Error setting argument 2");
+		ret_code = clSetKernelArg(nt_crypt_kernel, 3, sizeof(buffer_out) , (void*) &buffer_out);                       if_error_log (ret_code,"Error setting argument 3");
+		
+		MEM_FREE(source);
+	}
 #else
-	memset(nt_buffer1x,0,16*4*NT_NUM_KEYS);
+	static unsigned int last_i[NT_NUM_KEYS];
+	static char saved_plain[32*NT_NUM_KEYS];
+
+	static void fmt_NT_init(void)
+	{
+		memset(last_i,0,4*NT_NUM_KEYS);
+	#if defined(NT_X86_64)
+		memset(nt_buffer8x,0,16*4*NT_NUM_KEYS);
+	#elif defined(NT_SSE2)
+		memset(nt_buffer4x,0,64*4*NT_NUM_KEYS1);
+		memset(nt_buffer1x,0,16*4*NT_NUM_KEYS1);
+	#else
+		memset(nt_buffer1x,0,16*4*NT_NUM_KEYS);
+	#endif
+	}
 #endif
-}
 
 static char * nt_split(char *ciphertext, int index)
 {
@@ -266,7 +370,9 @@ static int binary_hash_2(void *binary)
 
 static int get_hash_0(int index)
 {
-#if defined(NT_X86_64)
+#ifdef NT_OPENCL
+	return bbbs[index] & 0x0F;
+#elif defined(NT_X86_64)
 	return output8x[32*(index>>3)+8+index%8] & 0x0F;
 #elif defined(NT_SSE2)
 	if(index<NT_NUM_KEYS4)
@@ -280,7 +386,9 @@ static int get_hash_0(int index)
 
 static int get_hash_1(int index)
 {
-#if defined(NT_X86_64)
+#ifdef NT_OPENCL
+	return bbbs[index] & 0xFF;
+#elif defined(NT_X86_64)
 	return output8x[32*(index>>3)+8+index%8] & 0xFF;
 #elif defined(NT_SSE2)
 	if(index<NT_NUM_KEYS4)
@@ -294,7 +402,9 @@ static int get_hash_1(int index)
 
 static int get_hash_2(int index)
 {
-#if defined(NT_X86_64)
+#ifdef NT_OPENCL
+	return bbbs[index] & 0x0FFF;
+#elif defined(NT_X86_64)
 	return output8x[32*(index>>3)+8+index%8] & 0x0FFF;
 #elif defined(NT_SSE2)
 	if(index<NT_NUM_KEYS4)
@@ -311,7 +421,11 @@ static int cmp_all(void *binary, int cou
 	unsigned int i=0;
 	unsigned int b=((unsigned int *)binary)[1];
 
-#if defined(NT_X86_64)
+#ifdef NT_OPENCL
+	for(;i<NT_NUM_KEYS;i++)
+		if(b==bbbs[i])
+			return 1;
+#elif defined(NT_X86_64)
 	for(;i<(NT_NUM_KEYS/2);i+=4)
 		if(b==output8x[i] || b==output8x[i+1] || b==output8x[i+2] || b==output8x[i+3] || b==output8x[i+4] || b==output8x[i+5] || b==output8x[i+6] || b==output8x[i+7])
 			return 1;
@@ -347,7 +461,23 @@ static int cmp_one(void * binary, int in
 	int pos2;
 	int pos3;
 	
-#if defined(NT_X86_64)
+#ifdef NT_OPENCL
+	cl_uint buffer_nt1[16];
+	cl_uint buffer_out1[4];
+	clEnqueueReadBuffer(queue, buffer_nt , CL_TRUE, 16*sizeof(cl_uint)*(index), sizeof(buffer_nt1 ), (void*)buffer_nt1 , 0, NULL, NULL);
+	clEnqueueReadBuffer(queue, buffer_out, CL_TRUE,  4*sizeof(cl_uint)*(index), sizeof(buffer_out1), (void*)buffer_out1, 0, NULL, NULL);
+	
+	buffer = (unsigned int *)buffer_nt1;
+	
+	a=buffer_out1[0];
+	b=buffer_out1[1];
+	c=buffer_out1[2];
+	d=buffer_out1[3];
+	
+	pos1=3;
+	pos2=(8+pos1);
+	pos3=(4+pos1);
+#elif defined(NT_X86_64)
 	int temp;
 	buffer=nt_buffer8x;
 	
@@ -431,54 +561,26 @@ static void set_salt(void *salt)
 {
 }
 
-static void set_key(char *key, int index)
-{
-	unsigned int i=0;
-	unsigned int md4_size=0;
-	unsigned int saved_base=index<<5;
-	unsigned int temp;
-	int buff_base;
-#if defined(NT_X86_64)
-	unsigned int last_length=last_i[index]<<2;
-
-	buff_base=128*(index>>3)+index%8;
-	
-	for(;key[md4_size] && md4_size<PLAINTEXT_LENGTH;i+=8,md4_size++)
+#ifdef NT_OPENCL
+	static void set_key(char *key, int index)
 	{
-		saved_plain[saved_base+md4_size]=key[md4_size];
-		temp=key[++md4_size];
-		saved_plain[saved_base+md4_size]=temp;
-		
-		if(temp)
-		{
-			nt_buffer8x[i+buff_base] = key[md4_size-1] | (temp<<16);
-		}
-		else
-		{
-			nt_buffer8x[i+buff_base] = key[md4_size-1] | 0x800000;
-			goto key_cleaning;
-		}
+		strcpy(saved_plain+(index<<5),key);
 	}
-	
-	nt_buffer8x[i+buff_base]=0x80;
-	saved_plain[saved_base+md4_size]=0;
-	
-key_cleaning:
-	i+=8;
-	for(;i<=last_length;i+=8)
-		nt_buffer8x[i+buff_base]=0;
-	
-	last_i[index]=md4_size;
-	
-	nt_buffer8x[112+buff_base] = md4_size << 4;
-#elif defined(NT_SSE2)
-	unsigned int last_length=last_i[index]<<1;
-
-	if(index<NT_NUM_KEYS4)
+#else
+	static void set_key(char *key, int index)
 	{
-		buff_base=64*(index>>2)+index%4;
+
+		unsigned int i=0;
+		unsigned int md4_size=0;
+		unsigned int saved_base=index<<5;
+		unsigned int temp;
+		int buff_base;
+	#if defined(NT_X86_64)
+		unsigned int last_length=last_i[index]<<2;
+
+		buff_base=128*(index>>3)+index%8;
 		
-		for(;key[md4_size] && md4_size<PLAINTEXT_LENGTH;i+=4,md4_size++)
+		for(;key[md4_size] && md4_size<PLAINTEXT_LENGTH;i+=8,md4_size++)
 		{
 			saved_plain[saved_base+md4_size]=key[md4_size];
 			temp=key[++md4_size];
@@ -486,30 +588,99 @@ key_cleaning:
 			
 			if(temp)
 			{
-				nt_buffer4x[i+buff_base] = key[md4_size-1] | (temp<<16);
+				nt_buffer8x[i+buff_base] = key[md4_size-1] | (temp<<16);
 			}
 			else
 			{
-				nt_buffer4x[i+buff_base] = key[md4_size-1] | 0x800000;
+				nt_buffer8x[i+buff_base] = key[md4_size-1] | 0x800000;
 				goto key_cleaning;
 			}
 		}
 		
-		nt_buffer4x[i+buff_base]=0x80;
+		nt_buffer8x[i+buff_base]=0x80;
 		saved_plain[saved_base+md4_size]=0;
 		
 	key_cleaning:
-		i+=4;
-		for(;i<=last_length;i+=4)
-			nt_buffer4x[i+buff_base]=0;
+		i+=8;
+		for(;i<=last_length;i+=8)
+			nt_buffer8x[i+buff_base]=0;
 		
 		last_i[index]=md4_size;
 		
-		nt_buffer4x[56+buff_base] = md4_size << 4;
-	}
-	else
-	{
-		buff_base=16*(index-NT_NUM_KEYS4);
+		nt_buffer8x[112+buff_base] = md4_size << 4;
+	#elif defined(NT_SSE2)
+		unsigned int last_length=last_i[index]<<1;
+
+		if(index<NT_NUM_KEYS4)
+		{
+			buff_base=64*(index>>2)+index%4;
+			
+			for(;key[md4_size] && md4_size<PLAINTEXT_LENGTH;i+=4,md4_size++)
+			{
+				saved_plain[saved_base+md4_size]=key[md4_size];
+				temp=key[++md4_size];
+				saved_plain[saved_base+md4_size]=temp;
+				
+				if(temp)
+				{
+					nt_buffer4x[i+buff_base] = key[md4_size-1] | (temp<<16);
+				}
+				else
+				{
+					nt_buffer4x[i+buff_base] = key[md4_size-1] | 0x800000;
+					goto key_cleaning;
+				}
+			}
+			
+			nt_buffer4x[i+buff_base]=0x80;
+			saved_plain[saved_base+md4_size]=0;
+			
+		key_cleaning:
+			i+=4;
+			for(;i<=last_length;i+=4)
+				nt_buffer4x[i+buff_base]=0;
+			
+			last_i[index]=md4_size;
+			
+			nt_buffer4x[56+buff_base] = md4_size << 4;
+		}
+		else
+		{
+			buff_base=16*(index-NT_NUM_KEYS4);
+			
+			for(;key[md4_size] && md4_size<PLAINTEXT_LENGTH;i++,md4_size++)
+			{
+				saved_plain[saved_base+md4_size]=key[md4_size];
+				temp=key[++md4_size];
+				saved_plain[saved_base+md4_size]=temp;
+				
+				if(temp)
+				{
+					nt_buffer1x[i+buff_base] = key[md4_size-1] | (temp<<16);
+				}
+				else
+				{
+					nt_buffer1x[i+buff_base] = key[md4_size-1] | 0x800000;
+					goto key_cleaning1;
+				}
+			}
+			
+			nt_buffer1x[i+buff_base]=0x80;
+			saved_plain[saved_base+md4_size]=0;
+			
+		key_cleaning1:
+			i++;
+			last_length>>=1;
+			
+			for(;i<=last_length;i++)
+				nt_buffer1x[i+buff_base]=0;
+			
+			last_i[index]=md4_size>>1;
+			
+			nt_buffer1x[14+buff_base] = md4_size << 4;
+		}
+	#else
+		buff_base=index<<4;
 		
 		for(;key[md4_size] && md4_size<PLAINTEXT_LENGTH;i++,md4_size++)
 		{
@@ -519,62 +690,29 @@ key_cleaning:
 			
 			if(temp)
 			{
-				nt_buffer1x[i+buff_base] = key[md4_size-1] | (temp<<16);
+				nt_buffer1x[buff_base+i] = key[md4_size-1] | (temp<<16);
 			}
 			else
 			{
-				nt_buffer1x[i+buff_base] = key[md4_size-1] | 0x800000;
-				goto key_cleaning1;
+				nt_buffer1x[buff_base+i] = key[md4_size-1] | 0x800000;
+				goto key_cleaning;
 			}
 		}
 		
-		nt_buffer1x[i+buff_base]=0x80;
+		nt_buffer1x[buff_base+i]=0x80;
 		saved_plain[saved_base+md4_size]=0;
 		
-	key_cleaning1:
+	key_cleaning:
 		i++;
-		last_length>>=1;
-		
-		for(;i<=last_length;i++)
-			nt_buffer1x[i+buff_base]=0;
+		for(;i<=last_i[index];i++)
+			nt_buffer1x[buff_base+i]=0;
 		
 		last_i[index]=md4_size>>1;
 		
-		nt_buffer1x[14+buff_base] = md4_size << 4;
+		nt_buffer1x[buff_base+14] = md4_size << 4;
+	#endif
 	}
-#else
-	buff_base=index<<4;
-	
-	for(;key[md4_size] && md4_size<PLAINTEXT_LENGTH;i++,md4_size++)
-	{
-		saved_plain[saved_base+md4_size]=key[md4_size];
-		temp=key[++md4_size];
-		saved_plain[saved_base+md4_size]=temp;
-		
-		if(temp)
-		{
-			nt_buffer1x[buff_base+i] = key[md4_size-1] | (temp<<16);
-		}
-		else
-		{
-			nt_buffer1x[buff_base+i] = key[md4_size-1] | 0x800000;
-			goto key_cleaning;
-		}
-	}
-	
-	nt_buffer1x[buff_base+i]=0x80;
-	saved_plain[saved_base+md4_size]=0;
-	
-key_cleaning:
-	i++;
-	for(;i<=last_i[index];i++)
-		nt_buffer1x[buff_base+i]=0;
-	
-	last_i[index]=md4_size>>1;
-	
-	nt_buffer1x[buff_base+14] = md4_size << 4;
 #endif
-}
 
 static char *get_key(int index)
 {