diff -urpN john-1.7.6.orig/README_CUDA john-1.7.6/README_CUDA
--- john-1.7.6.orig/README_CUDA	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.6/README_CUDA	2011-04-04 01:04:52.000000000 +0000
@@ -0,0 +1,12 @@
+README:
+SHA256CUDA is supported/tested only on following builds:
+linux-x86-mmx
+linux-x86-64
+linux-x86-any
+If you have problems while building patched JtR:
+  Please check that paths to yours CUDA library are proper in Makefile lines:
+    23: CUDAPATH = /usr/local/cuda/lib
+    24: CUDA64PATH = /usr/local/cuda/lib64
+  Please check that yours PATH contains cuda lib,bin and include paths.
+  Please check that yours LD_LIBRARY_PATH contains cuda lib path
+  You can contact me at lukas[dot]odzioba[at]gmail[dot]com
\ Brak znaku nowej linii na końcu pliku
Pliki binarne john-1.7.6.orig/run/john i john-1.7.6/run/john się różnią
Pliki binarne john-1.7.6.orig/run/unafs i john-1.7.6/run/unafs się różnią
Pliki binarne john-1.7.6.orig/run/unique i john-1.7.6/run/unique się różnią
Pliki binarne john-1.7.6.orig/run/unshadow i john-1.7.6/run/unshadow się różnią
Pliki binarne john-1.7.6.orig/src/AFS_fmt.o i john-1.7.6/src/AFS_fmt.o się różnią
diff -urpN john-1.7.6.orig/src/arch.h john-1.7.6/src/arch.h
--- john-1.7.6.orig/src/arch.h	2011-04-08 18:27:31.532394474 +0000
+++ john-1.7.6/src/arch.h	1970-01-01 00:00:00.000000000 +0000
@@ -1,70 +0,0 @@
-/*
- * This file is part of John the Ripper password cracker,
- * Copyright (c) 2003,2006,2008,2010 by Solar Designer
- */
-
-/*
- * Architecture specific parameters for x86-64.
- */
-
-#ifndef _JOHN_ARCH_H
-#define _JOHN_ARCH_H
-
-#define ARCH_WORD			long
-#define ARCH_SIZE			8
-#define ARCH_BITS			64
-#define ARCH_BITS_LOG			6
-#define ARCH_BITS_STR			"64"
-#define ARCH_LITTLE_ENDIAN		1
-#define ARCH_INT_GT_32			0
-#define ARCH_ALLOWS_UNALIGNED		1
-#define ARCH_INDEX(x)			((unsigned int)(unsigned char)(x))
-
-#define OS_TIMER			1
-#define OS_FLOCK			1
-
-#define CPU_DETECT			0
-
-#define DES_ASM				0
-#define DES_128K			0
-#define DES_X2				0
-#define DES_MASK			1
-#define DES_SCALE			1
-#define DES_EXTB			1
-#define DES_COPY			0
-#if defined(__SSE2__) && \
-    ((__GNUC__ == 4 && __GNUC_MINOR__ >= 4) || __GNUC__ > 4)
-#define DES_BS_ASM			0
-#if 1
-#define DES_BS_VECTOR			2
-#define DES_BS_ALGORITHM_NAME		"128/128 BS SSE2-16"
-#elif 0
-#define DES_BS_VECTOR			3
-#define DES_BS_VECTOR_SIZE		4
-#define DES_BS_ALGORITHM_NAME		"128/128 BS SSE2-16 + 64/64 BS MMX"
-#elif 0
-#define DES_BS_NO_MMX
-#define DES_BS_VECTOR			3
-#define DES_BS_VECTOR_SIZE		4
-#define DES_BS_ALGORITHM_NAME		"128/128 BS SSE2-16 + 64/64 BS"
-#else
-#define DES_BS_VECTOR			4
-#define DES_BS_ALGORITHM_NAME		"128/128 BS SSE2-16 + 64/64 BS MMX + 64/64 BS"
-#endif
-#else
-#define DES_BS_ASM			1
-#define DES_BS_VECTOR			2
-#define DES_BS_ALGORITHM_NAME		"128/128 BS SSE2-16"
-#endif
-#define DES_BS				1
-#define DES_BS_EXPAND			1
-
-#define MD5_ASM				0
-#define MD5_X2				1
-#define MD5_IMM				1
-
-#define BF_ASM				0
-#define BF_SCALE			1
-#define BF_X2				1
-
-#endif
Pliki binarne john-1.7.6.orig/src/batch.o i john-1.7.6/src/batch.o się różnią
Pliki binarne john-1.7.6.orig/src/bench.o i john-1.7.6/src/bench.o się różnią
Pliki binarne john-1.7.6.orig/src/BF_fmt.o i john-1.7.6/src/BF_fmt.o się różnią
Pliki binarne john-1.7.6.orig/src/BF_std.o i john-1.7.6/src/BF_std.o się różnią
Pliki binarne john-1.7.6.orig/src/BSDI_fmt.o i john-1.7.6/src/BSDI_fmt.o się różnią
Pliki binarne john-1.7.6.orig/src/c3_fmt.o i john-1.7.6/src/c3_fmt.o się różnią
Pliki binarne john-1.7.6.orig/src/charset.o i john-1.7.6/src/charset.o się różnią
Pliki binarne john-1.7.6.orig/src/common.o i john-1.7.6/src/common.o się różnią
Pliki binarne john-1.7.6.orig/src/compiler.o i john-1.7.6/src/compiler.o się różnią
Pliki binarne john-1.7.6.orig/src/config.o i john-1.7.6/src/config.o się różnią
Pliki binarne john-1.7.6.orig/src/cracker.o i john-1.7.6/src/cracker.o się różnią
Pliki binarne john-1.7.6.orig/src/crc32.o i john-1.7.6/src/crc32.o się różnią
diff -urpN john-1.7.6.orig/src/cuda/sha256.cu john-1.7.6/src/cuda/sha256.cu
--- john-1.7.6.orig/src/cuda/sha256.cu	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.6/src/cuda/sha256.cu	2011-04-14 21:24:01.702532500 +0000
@@ -0,0 +1,145 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include <assert.h>
+#include <string.h>
+#include "../sha256.h"
+
+const uint DATA_IN_SIZE=KEYS_PER_CRYPT*sizeof(sha256_password);
+const uint DATA_OUT_SIZE=KEYS_PER_CRYPT*sizeof(sha256_hash);
+
+__constant__ uint32_t _h[8];
+__constant__ uint32_t k[64];
+
+sha256_password *cuda_data0=NULL;		///candidates
+sha256_password *cuda_data1=NULL;		///candidates
+sha256_password *cuda_data2=NULL;		///candidates
+
+sha256_hash *cuda_data_out0=NULL;		///sha256(candidates)
+sha256_hash *cuda_data_out1=NULL;		///sha256(candidates)
+sha256_hash *cuda_data_out2=NULL;		///sha256(candidates)
+
+
+sha256_password *host_data=NULL;
+sha256_hash *host_data_out=NULL;
+
+cudaStream_t stream0,stream1,stream2;
+
+extern "C"  void mem_init(sha256_password *,sha256_hash*);
+extern "C"  void const_mem_cpy(void);
+extern "C"  void mem_clear(void);
+extern "C"  void sha256(void);
+extern "C" void* pageLockedMalloc(void *w,uint32_t size);
+extern "C"  void pageLockedFree(void *w);
+extern "C" void checkoverlapsupport(void);
+
+__global__ void kernel_sha256(sha256_password *data,sha256_hash* data_out);
+
+
+void * pageLockedMalloc(void *w,uint32_t size)
+{
+  cudaHostAlloc((void**)&w,size,cudaHostAllocDefault); 
+  return w;
+}
+void pageLockedFree(void *w)
+{
+  cudaFreeHost(w);
+}
+
+void checkoverlapsupport()
+{
+  cudaDeviceProp prop;
+  int dev;
+  cudaGetDevice(&dev);
+  cudaGetDeviceProperties(&prop,dev);
+  if(!prop.deviceOverlap){
+    printf("Overlaping is not supported by your device\n");
+  }
+}
+
+void sha256(void)
+{
+  dim3 dimGrid(BLOCKS/3);	/// todo auto detection 
+  dim3 dimBlock(THREADS);
+
+  cudaMemcpyAsync(cuda_data0,host_data,DATA_IN_SIZE/3,cudaMemcpyHostToDevice,stream0);
+  kernel_sha256<<<dimGrid,dimBlock,0,stream0>>>(cuda_data0,cuda_data_out0);
+  cudaMemcpyAsync(host_data_out,cuda_data_out0,DATA_OUT_SIZE/3,cudaMemcpyDeviceToHost,stream0);
+ 
+  cudaMemcpyAsync(cuda_data1,host_data+DATA_IN_SIZE/3,DATA_IN_SIZE/3,cudaMemcpyHostToDevice,stream1);
+  kernel_sha256<<<dimGrid,dimBlock,0,stream1>>>(cuda_data1,cuda_data_out1);
+  cudaMemcpyAsync(host_data_out+DATA_OUT_SIZE/3,cuda_data_out1,DATA_OUT_SIZE/3,cudaMemcpyDeviceToHost,stream1);
+
+  cudaMemcpyAsync(cuda_data1,host_data+2*DATA_IN_SIZE/3,DATA_IN_SIZE/3,cudaMemcpyHostToDevice,stream1);
+  kernel_sha256<<<dimGrid,dimBlock,0,stream2>>>(cuda_data1,cuda_data_out1);
+  cudaMemcpyAsync(host_data_out+2*DATA_OUT_SIZE/3,cuda_data_out1,DATA_OUT_SIZE/3,cudaMemcpyDeviceToHost,stream1);
+}
+
+void const_mem_cpy(void)
+{
+  cudaMemcpyToSymbol(_h,host_h,sizeof(uint32_t)*8);
+  cudaMemcpyToSymbol(k,host_k,sizeof(uint32_t)*64); 
+}
+
+void mem_init(sha256_password *p,sha256_hash *h)
+{
+  host_data=p;
+  host_data_out=h;
+  cudaMalloc(&cuda_data0,DATA_IN_SIZE/3);
+  cudaMalloc(&cuda_data1,DATA_IN_SIZE/3);
+  cudaMalloc(&cuda_data2,DATA_IN_SIZE/3);
+  cudaMalloc(&cuda_data_out0,DATA_OUT_SIZE/3);
+  cudaMalloc(&cuda_data_out1,DATA_OUT_SIZE/3);
+  cudaMalloc(&cuda_data_out2,DATA_OUT_SIZE/3);
+
+  cudaStreamCreate(&stream0);
+  cudaStreamCreate(&stream1);
+  cudaStreamCreate(&stream2);
+}
+
+void mem_clear()
+{
+  cudaStreamSynchronize(stream0);
+  cudaStreamSynchronize(stream1);
+  cudaStreamSynchronize(stream2);
+
+  cudaStreamDestroy(stream0);
+  cudaStreamDestroy(stream1);
+  cudaStreamDestroy(stream2);
+ 
+  cudaFree(cuda_data0);
+  cudaFree(cuda_data1);
+  cudaFree(cuda_data2);
+
+  cudaFree(cuda_data_out0);  
+  cudaFree(cuda_data_out1);  
+  cudaFree(cuda_data_out2); 
+}
+
+__global__ void kernel_sha256(sha256_password *data,sha256_hash* data_out){ 
+  uint32_t idx = blockIdx.x*blockDim.x + threadIdx.x;
+  
+  uint32_t w[64];
+  sha256_hash* out=&data_out[idx];
+
+  #pragma unroll 16
+  for(uint32_t j=0;j<16;j++)  w[j]=data[idx].v[j];
+  
+  #pragma unroll 48
+  for(uint32_t j=16;j<64;j++) w[j]=sigma1(w[j-2])+w[j-7]+sigma0(w[j-15])+w[j-16];
+  
+  
+  uint32_t a=_h[0];uint32_t b=_h[1];uint32_t c=_h[2];uint32_t d=_h[3];
+  uint32_t e=_h[4];uint32_t f=_h[5];uint32_t g=_h[6];uint32_t h=_h[7];
+  
+  uint32_t t1,t2;
+  #pragma unroll 64
+  for(uint32_t j=0;j<64;j++){	
+    t1=h+Sigma1(e)+Ch(e,f,g)+k[j]+w[j];
+    t2=Sigma0(a)+Maj(a,b,c);
+   h=g;g=f;f=e;e=d+t1;d=c;c=b;b=a;a=t1+t2;
+  }
+  out->v[0]=a+_h[0];out->v[1]=b+_h[1];
+  out->v[2]=c+_h[2];out->v[3]=d+_h[3];
+  out->v[4]=e+_h[4];out->v[5]=f+_h[5];
+  out->v[6]=g+_h[6];out->v[7]=h+_h[7];
+}
Pliki binarne john-1.7.6.orig/src/DES_bs_b.o i john-1.7.6/src/DES_bs_b.o się różnią
Pliki binarne john-1.7.6.orig/src/DES_bs.o i john-1.7.6/src/DES_bs.o się różnią
Pliki binarne john-1.7.6.orig/src/DES_fmt.o i john-1.7.6/src/DES_fmt.o się różnią
Pliki binarne john-1.7.6.orig/src/DES_std.o i john-1.7.6/src/DES_std.o się różnią
Pliki binarne john-1.7.6.orig/src/external.o i john-1.7.6/src/external.o się różnią
Pliki binarne john-1.7.6.orig/src/formats.o i john-1.7.6/src/formats.o się różnią
Pliki binarne john-1.7.6.orig/src/getopt.o i john-1.7.6/src/getopt.o się różnią
Pliki binarne john-1.7.6.orig/src/idle.o i john-1.7.6/src/idle.o się różnią
Pliki binarne john-1.7.6.orig/src/inc.o i john-1.7.6/src/inc.o się różnią
diff -urpN john-1.7.6.orig/src/john.c john-1.7.6/src/john.c
--- john-1.7.6.orig/src/john.c	2011-04-08 18:27:31.552382280 +0000
+++ john-1.7.6/src/john.c	2011-04-11 19:06:28.833460935 +0000
@@ -38,6 +38,7 @@ extern int CPU_detect(void);
 
 extern struct fmt_main fmt_DES, fmt_BSDI, fmt_MD5, fmt_BF;
 extern struct fmt_main fmt_AFS, fmt_LM;
+extern struct fmt_main fmt_SHA256CUDA;
 #ifdef HAVE_CRYPT
 extern struct fmt_main fmt_crypt;
 #endif
@@ -69,10 +70,12 @@ static void john_register_all(void)
 	john_register_one(&fmt_BF);
 	john_register_one(&fmt_AFS);
 	john_register_one(&fmt_LM);
+	john_register_one(&fmt_SHA256CUDA);
+
+
 #ifdef HAVE_CRYPT
 	john_register_one(&fmt_crypt);
 #endif
-
 	if (!fmt_list) {
 		fprintf(stderr, "Unknown ciphertext format name requested\n");
 		error();
Pliki binarne john-1.7.6.orig/src/john.o i john-1.7.6/src/john.o się różnią
Pliki binarne john-1.7.6.orig/src/list.o i john-1.7.6/src/list.o się różnią
Pliki binarne john-1.7.6.orig/src/LM_fmt.o i john-1.7.6/src/LM_fmt.o się różnią
Pliki binarne john-1.7.6.orig/src/loader.o i john-1.7.6/src/loader.o się różnią
Pliki binarne john-1.7.6.orig/src/logger.o i john-1.7.6/src/logger.o się różnią
diff -urpN john-1.7.6.orig/src/Makefile john-1.7.6/src/Makefile
--- john-1.7.6.orig/src/Makefile	2011-04-08 18:27:31.492383382 +0000
+++ john-1.7.6/src/Makefile	2011-04-11 19:22:03.233458102 +0000
@@ -2,7 +2,7 @@
 # This file is part of John the Ripper password cracker,
 # Copyright (c) 1996-2010 by Solar Designer
 #
-
+NVCC = nvcc
 CC = gcc
 AS = $(CC)
 LD = $(CC)
@@ -19,7 +19,11 @@ OMPFLAGS =
 #OMPFLAGS = -fopenmp
 # Sun Studio with OpenMP (set the OMP_NUM_THREADS env var at runtime)
 #OMPFLAGS = -xopenmp
-CFLAGS = -c -Wall -O2 -fomit-frame-pointer $(OMPFLAGS)
+NVCCFLAGS = -c --maxrregcount 14 -Xptxas -v
+CUDAPATH = /usr/local/cuda/lib
+CUDA64PATH = /usr/local/cuda/lib64
+CUDALIB=$(CUDAPATH)
+CFLAGS = -c -Wall -O2 -fomit-frame-pointer  $(OMPFLAGS)
 ASFLAGS = -c $(OMPFLAGS)
 LDFLAGS = -s $(OMPFLAGS)
 OPT_NORMAL = -funroll-loops
@@ -32,13 +36,14 @@ JOHN_OBJS = \
 	BF_fmt.o BF_std.o \
 	AFS_fmt.o \
 	LM_fmt.o \
+	sha256cuda_fmt.o sha256.o \
 	batch.o bench.o charset.o common.o compiler.o config.o cracker.o \
 	crc32.o external.o formats.o getopt.o idle.o inc.o john.o list.o \
 	loader.o logger.o math.o memory.o misc.o options.o params.o path.o \
 	recovery.o rpp.o rules.o signals.o single.o status.o tty.o wordlist.o \
 	unshadow.o \
 	unafs.o \
-	unique.o
+	unique.o 
 
 BENCH_DES_OBJS_ORIG = \
 	DES_fmt.o DES_std.o
@@ -148,7 +153,7 @@ linux-x86-64:
 	$(MAKE) $(PROJ) \
 		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86-64.o" \
 		CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \
-		LDFLAGS="$(LDFLAGS) -lcrypt"
+		LDFLAGS="$(LDFLAGS) -L$(CUDA64PATH) -lcudart -lcrypt" 
 
 linux-x86-64-32-sse2:
 	$(LN) x86-sse.h arch.h
@@ -178,14 +183,14 @@ linux-x86-mmx:
 	$(MAKE) $(PROJ) \
 		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86.o x86-mmx.o" \
 		CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \
-		LDFLAGS="$(LDFLAGS) -lcrypt"
+		LDFLAGS="$(LDFLAGS) -L$(CUDAPATH) -lcudart -lcrypt"
 
 linux-x86-any:
 	$(LN) x86-any.h arch.h
 	$(MAKE) $(PROJ) \
 		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86.o" \
 		CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \
-		LDFLAGS="$(LDFLAGS) -lcrypt"
+		LDFLAGS="$(LDFLAGS) -L$(CUDAPATH) -lcudart -lcrypt" 
 
 linux-alpha:
 	$(LN) alpha.h arch.h
@@ -776,6 +781,13 @@ generic.h:
 		"$(BENCH_MD5_OBJS_DEPEND)" \
 		"$(BENCH_BF_OBJS_DEPEND)"
 
+sha256.o:  sha256.h cuda/sha256.cu
+	$(shell  cd cuda; nvcc $(NVCCFLAGS) sha256.cu)
+
+sha256cuda_fmt.o: sha256.o sha256cuda_fmt.c sha256.h
+	cp cuda/sha256.o sha256.o
+	$(CC)  $(CFLAGS) sha256cuda_fmt.c
+
 bench: $(BENCH_OBJS)
 	$(LD) $(BENCH_OBJS) $(LDFLAGS) -o bench
 
@@ -855,6 +867,7 @@ depend:
 
 clean:
 	$(RM) $(PROJ) $(PROJ_DOS) $(PROJ_WIN32)
+	$(RM) cuda/sha256.o
 	$(RM) ../run/john.exe john-macosx-* *.o *.bak core
 	$(RM) detect bench generic.h arch.h tmp.s
 	$(CP) $(NULL) Makefile.dep
Pliki binarne john-1.7.6.orig/src/math.o i john-1.7.6/src/math.o się różnią
Pliki binarne john-1.7.6.orig/src/MD5_fmt.o i john-1.7.6/src/MD5_fmt.o się różnią
Pliki binarne john-1.7.6.orig/src/MD5_std.o i john-1.7.6/src/MD5_std.o się różnią
Pliki binarne john-1.7.6.orig/src/memory.o i john-1.7.6/src/memory.o się różnią
Pliki binarne john-1.7.6.orig/src/misc.o i john-1.7.6/src/misc.o się różnią
diff -urpN john-1.7.6.orig/src/options.c john-1.7.6/src/options.c
--- john-1.7.6.orig/src/options.c	2011-04-08 18:27:31.562365113 +0000
+++ john-1.7.6/src/options.c	2011-04-11 19:06:28.833460935 +0000
@@ -101,7 +101,7 @@ static struct opt_entry opt_list[] = {
 "--salts=[-]COUNT           load salts with[out] at least COUNT passwords " \
 	"only\n" \
 "--format=NAME              force hash type NAME: " \
-	"DES/BSDI/MD5/BF/AFS/LM/crypt\n" \
+	"DES/BSDI/MD5/BF/AFS/LM/crypt/SHA256CUDA\n" \
 "--save-memory=LEVEL        enable memory saving, at LEVEL 1..3\n"
 
 void opt_init(char *name, int argc, char **argv)
Pliki binarne john-1.7.6.orig/src/options.o i john-1.7.6/src/options.o się różnią
Pliki binarne john-1.7.6.orig/src/params.o i john-1.7.6/src/params.o się różnią
Pliki binarne john-1.7.6.orig/src/path.o i john-1.7.6/src/path.o się różnią
Pliki binarne john-1.7.6.orig/src/recovery.o i john-1.7.6/src/recovery.o się różnią
Pliki binarne john-1.7.6.orig/src/rpp.o i john-1.7.6/src/rpp.o się różnią
Pliki binarne john-1.7.6.orig/src/rules.o i john-1.7.6/src/rules.o się różnią
diff -urpN john-1.7.6.orig/src/sha256cuda_fmt.c john-1.7.6/src/sha256cuda_fmt.c
--- john-1.7.6.orig/src/sha256cuda_fmt.c	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.6/src/sha256cuda_fmt.c	2011-04-12 18:27:54.871569863 +0000
@@ -0,0 +1,233 @@
+/*
+* This software is Copyright (c) 2011 Lukas Odzioba
+* <lukas dot odzioba at gmail dot com> 
+* and it is hereby released to the general public under the following terms:
+* Redistribution and use in source and binary forms, with or without modification, are permitted.
+*/
+#include <string.h>
+#include "arch.h"
+#include "formats.h"
+#include "common.h"	///for ARCH_WORD_32,atoi16
+#include "misc.h"	///for strnzcpy
+#include "sha256.h"
+
+#define FORMAT_LABEL		"sha256cuda" 
+#define FORMAT_NAME		"SHA256CUDA"
+
+#define SHA_TYPE		"SHA256" 
+
+#define BENCHMARK_COMMENT	""
+#define BENCHMARK_LENGTH	-1 /// Raw benchmark
+
+#define PLAINTEXT_LENGTH	54
+#define CIPHERTEXT_LENGTH	64 ///256bit
+#define BINARY_SIZE		32 
+
+#define SALT_SIZE		0
+
+#define MIN_KEYS_PER_CRYPT	KEYS_PER_CRYPT	
+#define MAX_KEYS_PER_CRYPT	KEYS_PER_CRYPT
+
+extern void mem_init(sha256_password *,sha256_hash*);
+extern void mem_clear(void);
+extern void sha256(void);
+extern void const_mem_cpy(void);
+extern void* pageLockedMalloc(void *w,uint32_t size);
+//extern void pageLockedFree(void *w);
+extern void checkoverlapsupport(void);
+
+static char saved_keys[MAX_KEYS_PER_CRYPT][PLAINTEXT_LENGTH+1];		/** plaintext ciphertexts **/
+static sha256_password 	*inbuffer;					/** binary ciphertexts **/
+static sha256_hash	*outbuffer;					/** calculated hashes **/
+
+static void preproc(char *key, int index){ /// todo - clean it
+  uint32_t dl=strlen(key),j;
+  uint32_t *blocks = inbuffer[index].v;
+  memset(inbuffer[index].v,0,sizeof(sha256_password));
+  for(j=0;j<dl;j++){
+      uint32_t tmp=0;
+      if(j%4==0) tmp|=(((uint32_t)key[j])<<24);
+      if(j%4==1) tmp|=(((uint32_t)key[j])<<16);
+      if(j%4==2) tmp|=(((uint32_t)key[j])<<8);
+      if(j%4==3) tmp|=((uint32_t)key[j]);
+      blocks[j/4]|=tmp;
+    }
+    if(j%4==0) blocks[j/4]|=(((uint32_t)0x80)<<24);
+    if(j%4==1) blocks[j/4]|=(((uint32_t)0x80)<<16);
+    if(j%4==2) blocks[j/4]|=(((uint32_t)0x80)<<8);
+    if(j%4==3) blocks[j/4]|=(((uint32_t)0x80));
+    blocks[15]=0x00000000|(dl*8);
+}
+
+static struct fmt_tests tests[]={
+  {"ba7816bf8f01cfea414140de5dae2223b00361a396177a9cb410ff61f20015ad","abc"},
+  {"96d9632f363564cc3032521409cf22a852f2032eec099ed5967c0d000cec607a","john"},
+  {"5d5b09f6dcb2d53a5fffc60c4ac0d55fabdf556069d6631545f42aa6e3500f2e","sha256"},
+  {"a49c2c9d0c006c8cb55a9a7a38822b83e0cd442614cb416af952fa50156761dc","openwall"},
+  {"e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855",""},
+  {NULL}  
+};
+
+static void init(void){
+  /*constant memory copy*/
+  const_mem_cpy();
+  uint32_t insize=sizeof(sha256_password)*MAX_KEYS_PER_CRYPT;
+  uint32_t outsize=sizeof(sha256_hash)*MAX_KEYS_PER_CRYPT;
+  inbuffer=pageLockedMalloc(inbuffer,insize);
+  outbuffer=pageLockedMalloc(outbuffer,outsize);
+  if(inbuffer==NULL||outbuffer==NULL){
+   printf("Page Locked memory alocation error");
+  }
+  checkoverlapsupport();
+}
+
+static int valid(char * ciphertext){
+  int i;
+  if(strlen(ciphertext)!=CIPHERTEXT_LENGTH) return 0;
+  for(i=0;i<CIPHERTEXT_LENGTH;i++){
+    if(!(
+      (ciphertext[i]>='0' && ciphertext[i]<='9')||
+      (ciphertext[i]>='a' && ciphertext[i]<='f')||
+      (ciphertext[i]>='A' && ciphertext[i]<='Z')))
+	return 0;
+  }
+  return 1;
+};
+
+
+static void *binary(char *ciphertext){
+  static char realcipher[BINARY_SIZE];
+  memset(realcipher,0,BINARY_SIZE);
+  int i;
+  for(i=0;i<BINARY_SIZE;i+=4){
+      realcipher[i]=atoi16[ARCH_INDEX(ciphertext[(i+3)*2])]*16+atoi16[ARCH_INDEX(ciphertext[(i+3)*2+1])];
+      realcipher[i+1]=atoi16[ARCH_INDEX(ciphertext[(i+2)*2])]*16+atoi16[ARCH_INDEX(ciphertext[(i+2)*2+1])];
+      realcipher[i+2]=atoi16[ARCH_INDEX(ciphertext[(i+1)*2])]*16+atoi16[ARCH_INDEX(ciphertext[(i+1)*2+1])];
+      realcipher[i+3]=atoi16[ARCH_INDEX(ciphertext[(i)*2])]*16+atoi16[ARCH_INDEX(ciphertext[(i)*2+1])];
+  }
+  return (void*)realcipher;
+}
+
+static int binary_hash_0(void *binary){
+  return (((ARCH_WORD_32*)binary)[0] & 0xf);
+}
+
+static int binary_hash_1(void *binary){
+  return ((ARCH_WORD_32*)binary)[0] & 0xff;
+}
+
+static int binary_hash_2(void *binary){
+  return ((ARCH_WORD_32*)binary)[0] & 0xfff;
+}
+
+static int binary_hash_3(void *binary){
+  return ((ARCH_WORD_32*)binary)[0] & 0xffff;
+}
+
+static int binary_hash_4(void *binary){
+  return ((ARCH_WORD_32*)binary)[0] & 0xfffff;
+}
+
+static void set_salt(void *salt){}
+static void set_key(char *key, int index){
+    memset(saved_keys[index],0,PLAINTEXT_LENGTH+1);
+    strnzcpy(saved_keys[index],key,PLAINTEXT_LENGTH);
+    preproc(key,index); ///todo - check that it is faster to do it on gpu
+}
+static char *get_key(int index){
+  return saved_keys[index];
+}
+
+static void crypt_all(int count){
+  mem_init(inbuffer,outbuffer);
+  sha256();
+  mem_clear(); ///hurts for small number KEYS_PER_CRYPT
+}
+
+static int get_hash_0(int index){
+  return ((ARCH_WORD_32*)outbuffer[index].v)[0] & 0xf;
+}
+
+static int get_hash_1(int index){
+  return ((ARCH_WORD_32*)outbuffer[index].v)[0] & 0xff;
+}
+
+static int get_hash_2(int index){
+  return ((ARCH_WORD_32*)outbuffer[index].v)[0] & 0xfff;
+}
+
+static int get_hash_3(int index){
+  return ((ARCH_WORD_32*)outbuffer[index].v)[0] & 0xffff;
+}
+
+static int get_hash_4(int index){
+  return ((ARCH_WORD_32*)outbuffer[index].v)[0] & 0xfffff;
+}
+
+static int cmp_all(void *binary,int count){
+  uint32_t i;
+  uint32_t 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){
+  int i;
+  uint32_t *t=(uint32_t *)binary;
+  for(i=0;i<8;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_SHA256CUDA={
+  {
+    FORMAT_LABEL,
+    FORMAT_NAME,
+    SHA_TYPE,
+    BENCHMARK_COMMENT,
+    BENCHMARK_LENGTH,
+    PLAINTEXT_LENGTH,
+    BINARY_SIZE,
+    SALT_SIZE,
+    MIN_KEYS_PER_CRYPT,
+    MAX_KEYS_PER_CRYPT,
+    FMT_CASE | FMT_8_BIT ,
+    tests
+  },
+  {
+    init,
+    valid,
+    fmt_default_split,
+    binary,
+    fmt_default_salt,
+    {
+      binary_hash_0,
+      binary_hash_1,
+      binary_hash_2,
+      binary_hash_3,
+      binary_hash_4
+    },
+    fmt_default_salt_hash,
+    set_salt,
+    set_key,
+    get_key,
+    fmt_default_clear_keys,
+    crypt_all,
+    {
+      get_hash_0,
+      get_hash_1,
+      get_hash_2,
+      get_hash_3,
+      get_hash_4
+    },
+    cmp_all,
+    cmp_one,
+    cmp_exact
+  }
+};
diff -urpN john-1.7.6.orig/src/sha256.h john-1.7.6/src/sha256.h
--- john-1.7.6.orig/src/sha256.h	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.6/src/sha256.h	2011-04-14 20:36:51.872532321 +0000
@@ -0,0 +1,42 @@
+#ifndef _SHA256_H
+#define _SHA256_H
+
+#ifndef uint32_t
+  #define uint32_t unsigned int  
+#endif
+
+#define rol(x,n) ((x << n) | (x >> (32-n)))	
+#define ror(x,n) ((x >> n) | (x << (32-n))) 	
+#define Ch(x,y,z) ((x & y) ^ ( (~x) & z)) 
+#define Maj(x,y,z) ((x & y) ^ (x & z) ^ (y & z))
+#define Sigma0(x) ((ror(x,2))  ^ (ror(x,13)) ^ (ror(x,22)))
+#define Sigma1(x) ((ror(x,6))  ^ (ror(x,11)) ^ (ror(x,25)))
+#define sigma0(x) ((ror(x,7))  ^ (ror(x,18)) ^(x>>3))
+#define sigma1(x) ((ror(x,17)) ^ (ror(x,19)) ^(x>>10))
+
+#define THREADS 128
+#define BLOCKS 4096*3
+#define KEYS_PER_CRYPT THREADS*BLOCKS
+
+typedef struct{
+  uint32_t v[16];  				///512bit 
+}sha256_password;
+
+typedef struct{
+  uint32_t v[8]; 				///256bit
+}sha256_hash;
+
+const uint32_t host_h[]={
+   0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0xa54ff53a, 0x510e527f, 0x9b05688c, 0x1f83d9ab, 0x5be0cd19};
+const uint32_t host_k[]={
+   0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, 0x923f82a4, 0xab1c5ed5,
+   0xd807aa98, 0x12835b01, 0x243185be, 0x550c7dc3, 0x72be5d74, 0x80deb1fe, 0x9bdc06a7, 0xc19bf174,
+   0xe49b69c1, 0xefbe4786, 0x0fc19dc6, 0x240ca1cc, 0x2de92c6f, 0x4a7484aa, 0x5cb0a9dc, 0x76f988da,
+   0x983e5152, 0xa831c66d, 0xb00327c8, 0xbf597fc7, 0xc6e00bf3, 0xd5a79147, 0x06ca6351, 0x14292967,
+   0x27b70a85, 0x2e1b2138, 0x4d2c6dfc, 0x53380d13, 0x650a7354, 0x766a0abb, 0x81c2c92e, 0x92722c85,
+   0xa2bfe8a1, 0xa81a664b, 0xc24b8b70, 0xc76c51a3, 0xd192e819, 0xd6990624, 0xf40e3585, 0x106aa070,
+   0x19a4c116, 0x1e376c08, 0x2748774c, 0x34b0bcb5, 0x391c0cb3, 0x4ed8aa4a, 0x5b9cca4f, 0x682e6ff3,
+   0x748f82ee, 0x78a5636f, 0x84c87814, 0x8cc70208, 0x90befffa, 0xa4506ceb, 0xbef9a3f7, 0xc67178f2};
+ 
+
+#endif
Pliki binarne john-1.7.6.orig/src/signals.o i john-1.7.6/src/signals.o się różnią
Pliki binarne john-1.7.6.orig/src/single.o i john-1.7.6/src/single.o się różnią
Pliki binarne john-1.7.6.orig/src/status.o i john-1.7.6/src/status.o się różnią
Pliki binarne john-1.7.6.orig/src/tty.o i john-1.7.6/src/tty.o się różnią
Pliki binarne john-1.7.6.orig/src/unafs.o i john-1.7.6/src/unafs.o się różnią
Pliki binarne john-1.7.6.orig/src/unique.o i john-1.7.6/src/unique.o się różnią
Pliki binarne john-1.7.6.orig/src/unshadow.o i john-1.7.6/src/unshadow.o się różnią
Pliki binarne john-1.7.6.orig/src/wordlist.o i john-1.7.6/src/wordlist.o się różnią
Pliki binarne john-1.7.6.orig/src/x86-64.o i john-1.7.6/src/x86-64.o się różnią
