diff -urpN john-1.7.7.orig//doc/CRYPTSHA256_CUDA john-1.7.7-crypt-sha256_8//doc/CRYPTSHA256_CUDA
--- john-1.7.7.orig//doc/CRYPTSHA256_CUDA	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.7-crypt-sha256_8//doc/CRYPTSHA256_CUDA	2011-07-06 10:47:08.146073714 +0000
@@ -0,0 +1,68 @@
+  ,-*
+ (_)   Release Notes
+
+==== Version        ====
+    CRYPTSHA256 CUDA ver. 2.0 patch for John The Ripper 1.7.7
+
+==== Release Date   ====
+    DD.MM.RRRR = 06.07.2011
+
+==== Authors        ====
+    Lukas Odzioba, lukas[dot]odzioba[at]gmail[dot]com
+
+==== Change Log     ====
+      ver 2:
+	Performance changes
+	Removed FULL_UNROLL option
+      ver 1:
+	Reimplemented from scratch
+      ver 0:
+	First revision
+       
+==== Restrictions   ====
+    Patch was tested/optimized for following configuration:
+    i3-2100, 4GB 1333Mhz, Slackware 13.1, GTX 460
+    Patch is not setting up THREADS and BLOCKS automaticaly, you should experiment with those constants 
+    to get better performance on your card.
+    Passwords length should be less than 16.
+
+==== ToDo 	    ====
+    Speed improvements
+
+==== Known Bugs     ====
+
+==== Environment    ====
+    CRYPTMD5CUDA 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:
+	24: CUDAPATH = /usr/local/cuda/lib
+	25: 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
+
+==== Instalation    ====
+    Apply this patch on top of JtR 1.7.7
+    $ cd john-1.7.7
+    $ patch -p1 < ../john-1.7.7-cryptsha256cuda-2.diff
+    $ cd john-1.7.7/src
+    for 32-bit systems:
+    $ make clean && make linux-x86-mmx
+    for 64-bit systems:
+    $ make clean && make linux-x86-64
+
+==== Usage          ====
+    Benchmark:
+    $ john-1.7.7/run/john -test --format=cryptsha256cuda
+
+    Incremental attack:
+    $ john-1.7.7/run/john -i -format=cryptsha256cuda FILE
+
+==== Benchmarks     ==== 
+    ver 2:
+      GeForce GTX 460 1024MB		Raw:    7798 c/s real, 7798 c/s virtual
+    ver 1:
+      GeForce GTX 460 1024MB		Raw:    6813 c/s real, 6813 c/s virtual
+    ver 0:
+      GeForce GTX 460 1024MB   		Raw:	4553 c/s real, 4553 c/s virtual
+
diff -urpN john-1.7.7.orig//src/Makefile john-1.7.7-crypt-sha256_8//src/Makefile
--- john-1.7.7.orig//src/Makefile	2011-02-27 12:31:50.000000000 +0000
+++ john-1.7.7-crypt-sha256_8//src/Makefile	2011-07-03 22:46:00.843525571 +0000
@@ -3,6 +3,7 @@
 # Copyright (c) 1996-2011 by Solar Designer
 #
 
+NVCC = nvcc
 CC = gcc
 AS = $(CC)
 LD = $(CC)
@@ -19,6 +20,10 @@ OMPFLAGS =
 #OMPFLAGS = -fopenmp
 # Sun Studio with OpenMP (set the OMP_NUM_THREADS env var at runtime)
 #OMPFLAGS = -xopenmp
+NVCCFLAGS = -c -Xptxas -v -arch=sm_10
+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)
@@ -33,6 +38,7 @@ JOHN_OBJS = \
 	AFS_fmt.o \
 	LM_fmt.o \
 	dummy.o \
+	cryptsha256cuda_fmt.o cryptsha256.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 \
@@ -153,14 +159,14 @@ 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-avx:
 	$(LN) x86-64.h arch.h
 	$(MAKE) $(PROJ) \
 		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o" \
 		CFLAGS="$(CFLAGS) -mavx -DHAVE_CRYPT" \
-		LDFLAGS="$(LDFLAGS) -lcrypt"
+		LDFLAGS="$(LDFLAGS)  -lcrypt"
 
 linux-x86-64-xop:
 	$(LN) x86-64.h arch.h
@@ -197,14 +203,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-x86-avx:
 	$(LN) x86-sse.h arch.h
@@ -809,6 +815,13 @@ generic.h:
 		"$(BENCH_MD5_OBJS_DEPEND)" \
 		"$(BENCH_BF_OBJS_DEPEND)"
 
+cryptsha256.o:  cryptsha256.h cuda/cryptsha256.cu
+	$(shell  cd cuda; nvcc $(NVCCFLAGS) cryptsha256.cu)
+	cp cuda/cryptsha256.o cryptsha256.o
+
+cryptsha256_fmt.o: cryptsha256.o cryptsha256cuda_fmt.c cryptsha256.h
+	$(CC)  $(CFLAGS) cryptsha256cuda_fmt.c
+
 bench: $(BENCH_OBJS)
 	$(LD) $(BENCH_OBJS) $(LDFLAGS) -o bench
 
@@ -888,6 +901,8 @@ depend:
 
 clean:
 	$(RM) $(PROJ) $(PROJ_DOS) $(PROJ_WIN32)
+	$(RM) cuda/cryptsha256.o cuda/*.c~
+	$(RM) *.cubin *.c~ *~
 	$(RM) ../run/john.exe john-macosx-* *.o *.bak core
 	$(RM) detect bench generic.h arch.h tmp.s
 	$(CP) $(NULL) Makefile.dep
diff -urpN john-1.7.7.orig//src/cryptsha256.h john-1.7.7-crypt-sha256_8//src/cryptsha256.h
--- john-1.7.7.orig//src/cryptsha256.h	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.7-crypt-sha256_8//src/cryptsha256.h	2011-07-06 10:48:45.916958081 +0000
@@ -0,0 +1,67 @@
+/*
+* This software is Copyright (c) 2011 Lukas Odzioba
+* <lukas dot odzioba at gmail dot com>
+* and it is hereby released to the general public under the following terms:
+* Redistribution and use in source and binary forms, with or without modification, are permitted.
+*/
+#ifndef _CRYPTSHA256_H
+#define _CRYPTSHA256_H
+
+#include <assert.h>
+#include "common.h"
+#include <stdint.h>
+#include <stdbool.h>
+
+#define BLOCKS 14
+#define THREADS 320
+
+#define KEYS_PER_CRYPT BLOCKS*THREADS
+
+#define uint32_t unsigned int
+#define uint8_t unsigned char
+
+#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 MAX(x,y) (x) > (y) ? (x) : (y)
+#define MIN(x,y) (x) < (y) ? (x) : (y)
+#define SHOW(x) printf("%s = %08x\n",#x,(x))
+# define SWAP(n) \
+    (((n) << 24) | (((n) & 0xff00) << 8) | (((n) >> 8) & 0xff00) | ((n) >> 24))
+
+#define SALT_LEN_MAX 16
+#define ROUNDS_DEFAULT 5000
+#define ROUNDS_MIN 1000
+#define ROUNDS_MAX 999999999
+
+typedef struct {
+	uint32_t H[8];
+	uint32_t total;
+	uint32_t buflen;
+	uint8_t buffer[64];
+} sha256_ctx;
+
+typedef struct {
+	unsigned char saltlen;
+	uint32_t rounds;
+	char salt[16];
+} crypt_sha256_salt;
+
+typedef struct {
+	unsigned char length;
+	unsigned char v[15];
+} crypt_sha256_password;
+
+
+#define hash_addr(j,idx) (((j)*(KEYS_PER_CRYPT))+(idx))
+typedef struct {
+	uint32_t v[8];		//256 bits
+} crypt_sha256_hash;
+
+#endif
diff -urpN john-1.7.7.orig//src/cryptsha256cuda_fmt.c john-1.7.7-crypt-sha256_8//src/cryptsha256cuda_fmt.c
--- john-1.7.7.orig//src/cryptsha256cuda_fmt.c	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.7-crypt-sha256_8//src/cryptsha256cuda_fmt.c	2011-07-06 10:48:54.085082773 +0000
@@ -0,0 +1,327 @@
+/*
+* 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"
+#include "misc.h"
+#include "cryptsha256.h"
+#include <unistd.h>
+
+#define FORMAT_LABEL		"cryptsha256cuda"
+#define FORMAT_NAME		"CRYPTSHA256CUDA"
+
+#define CRYPT_TYPE		"SHA256-based CRYPT"
+
+#define BENCHMARK_COMMENT	""
+#define BENCHMARK_LENGTH	-1
+
+#define PLAINTEXT_LENGTH	15
+#define CIPHERTEXT_LENGTH	34
+#define BINARY_SIZE		32
+#define MD5_DIGEST_LENGTH 	16
+
+#define SALT_SIZE		16
+
+#define MIN_KEYS_PER_CRYPT	KEYS_PER_CRYPT
+#define MAX_KEYS_PER_CRYPT	KEYS_PER_CRYPT
+
+extern void sha256_crypt_gpu(crypt_sha256_password * inbuffer,
+    uint32_t * outbuffer, crypt_sha256_salt * host_salt);
+
+static crypt_sha256_password inbuffer[MAX_KEYS_PER_CRYPT];			/** plaintext ciphertexts **/
+static uint32_t outbuffer[MAX_KEYS_PER_CRYPT * 8];				/** calculated hashes **/
+
+static char currentsalt[64];
+crypt_sha256_salt host_salt;
+
+void sha256_crypt_cpu(crypt_sha256_password * passwords,
+    crypt_sha256_hash * output, crypt_sha256_salt * salt);
+
+
+static struct fmt_tests tests[] = {
+	{"$5$saltstring$5B8vYYiY.CVt1RlTTf8KbXBH3hsxY/GNooZaBBGWEc5",
+	    "Hello world!"},
+	//{"$5$saltstring$0Az3qME7zTXm78kfHrR2OtT8WOu2gd8bcVn/9Y.3l/7", "john"},
+
+	{NULL}
+};
+
+static void init(void)
+{
+}
+
+static int valid(char *ciphertext)
+{
+	uint32_t i, j;
+	int len = strlen(ciphertext);
+
+	if (strncmp(ciphertext, "$5$", 3) != 0)
+		return 0;
+	char *p = strrchr(ciphertext, '$');
+	for (i = p - ciphertext + 1; i < len; i++) {
+		int found = 0;
+		for (j = 0; j < 64; j++)
+			if (itoa64[j] == ARCH_INDEX(ciphertext[i])) {
+				found = 1;
+				break;
+			}
+		if (found == 0)
+			return 0;
+	}
+	if (len - (p - ciphertext + 1) != 43)
+		return 0;
+	return 1;
+};
+
+int findb64(char c)
+{
+	int i = 0;
+	for (i = 0; i < 64; i++)
+		if (c == itoa64[i])
+			return i;
+	return 0;
+}
+
+void magic(char *crypt, char *alt)
+{
+
+#define _24bit_from_b64(I,B2,B1,B0) \
+  do{\
+      unsigned char c1=findb64(crypt[I+0]);\
+      unsigned char c2=findb64(crypt[I+1]);\
+      unsigned char c3=findb64(crypt[I+2]);\
+      unsigned char c4=findb64(crypt[I+3]);\
+      unsigned int w=c4<<18|c3<<12|c2<<6|c1;\
+      unsigned char b2=w&0xff;w>>=8;\
+      unsigned char b1=w&0xff;w>>=8;\
+      unsigned char b0=w&0xff;w>>=8;\
+      alt[B2]=b0;\
+      alt[B1]=b1;\
+      alt[B0]=b2;\
+  }while(0)
+
+	_24bit_from_b64(0, 0, 10, 20);
+	_24bit_from_b64(4, 21, 1, 11);
+	_24bit_from_b64(8, 12, 22, 2);
+	_24bit_from_b64(12, 3, 13, 23);
+	_24bit_from_b64(16, 24, 4, 14);
+	_24bit_from_b64(20, 15, 25, 5);
+	_24bit_from_b64(24, 6, 16, 26);
+	_24bit_from_b64(28, 27, 7, 17);
+	_24bit_from_b64(32, 18, 28, 8);
+	_24bit_from_b64(36, 9, 19, 29);
+	unsigned int w =
+	    findb64(crypt[42]) << 12 | findb64(crypt[41]) << 6 |
+	    findb64(crypt[40]);
+	alt[30] = w & 0xff;
+	w >>= 8;
+	alt[31] = w & 0xff;
+	w >>= 8;
+}
+
+
+static void *binary(char *ciphertext)
+{
+	static char b[BINARY_SIZE];
+	memset(b, 0, BINARY_SIZE);
+	char *p = strrchr(ciphertext, '$') + 1;
+	magic(p, b);
+	return (void *) b;
+}
+
+
+static void *salt(char *ciphertext)
+{
+	int end = 0, i, len = strlen(ciphertext);
+	for (i = len - 1; i >= 0; i--)
+		if (ciphertext[i] == '$') {
+			end = i;
+			break;
+
+		}
+
+	static unsigned char ret[50];
+	for (i = 0; i < end; i++)
+		ret[i] = ciphertext[i];
+	ret[end] = 0;
+	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 void set_salt(void *salt)
+{
+	unsigned char *s = salt;
+	int i, len = strlen(salt);
+	for (i = 0; i < len; i++)
+		currentsalt[i] = s[i];
+	currentsalt[len] = 0;
+	unsigned char offset = 0;
+	host_salt.rounds = ROUNDS_DEFAULT;
+
+	if (strncmp((char *) "$5$", (char *) currentsalt, 3) == 0)
+		offset += 3;
+
+	if (strncmp((char *) currentsalt + offset, (char *) "rounds=", 7) == 0) {
+		const char *num = currentsalt + offset + 7;
+		char *endp;
+		unsigned long int srounds = strtoul(num, &endp, 10);
+
+		if (*endp == '$') {
+			endp += 1;
+			host_salt.rounds =
+			    MAX(ROUNDS_MIN, MIN(srounds, ROUNDS_MAX));
+		}
+		offset = endp - currentsalt;
+	}
+	memcpy(host_salt.salt, currentsalt + offset, 16);
+	host_salt.saltlen = strlen(host_salt.salt);
+}
+
+static void set_key(char *key, int index)
+{
+	memset(inbuffer[index].v, 0, PLAINTEXT_LENGTH + 1);
+	int len = strlen(key);
+	inbuffer[index].length = len;
+	strnzcpy((char *) inbuffer[index].v, key, PLAINTEXT_LENGTH);
+}
+
+static char *get_key(int index)
+{
+	return (char *) inbuffer[index].v;
+}
+
+static void crypt_all(int count)
+{
+	sha256_crypt_gpu(inbuffer, outbuffer, &host_salt);
+}
+
+static int get_hash_0(int index)
+{
+
+	uint32_t *out = outbuffer;
+	return out[hash_addr(0, index)] & 0xf;
+}
+
+static int get_hash_1(int index)
+{
+	uint32_t *out = outbuffer;
+	return out[hash_addr(0, index)] & 0xff;
+}
+
+static int get_hash_2(int index)
+{
+	uint32_t *out = outbuffer;
+	return out[hash_addr(0, index)] & 0xfff;
+}
+
+static int get_hash_3(int index)
+{
+	uint32_t *out = outbuffer;
+	return out[hash_addr(0, index)] & 0xffff;
+}
+
+static int get_hash_4(int index)
+{
+	uint32_t *out = outbuffer;
+	return out[hash_addr(0, index)] & 0xfffff;
+}
+
+static int cmp_all(void *binary, int count)
+{
+	uint32_t i;
+	uint32_t b = ((uint32_t *) binary)[0];
+	uint32_t *out = outbuffer;
+	for (i = 0; i < count; i++)
+		if (b == out[hash_addr(0, i)])
+			return 1;
+	return 0;
+}
+
+static int cmp_one(void *binary, int index)
+{
+	int i;
+	uint32_t *t = (uint32_t *) binary;
+	uint32_t *out = outbuffer;
+
+	for (i = 0; i < 8; i++)
+		if (t[i] != out[hash_addr(i, index)])
+			return 0;
+	return 1;
+}
+
+static int cmp_exact(char *source, int count)
+{
+	return 1;
+}
+
+struct fmt_main fmt_CRYPTSHA256CUDA = {
+	{
+		    FORMAT_LABEL,
+		    FORMAT_NAME,
+		    CRYPT_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,
+		    valid,
+		    fmt_default_split,
+		    binary,
+		    salt,
+		    {
+				binary_hash_0,
+				binary_hash_1,
+				binary_hash_2,
+				binary_hash_3,
+			binary_hash_4},
+		    fmt_default_salt_hash,
+		    set_salt,
+		    set_key,
+		    get_key,
+		    fmt_default_clear_keys,
+		    crypt_all,
+		    {
+				get_hash_0,
+				get_hash_1,
+				get_hash_2,
+				get_hash_3,
+			get_hash_4},
+		    cmp_all,
+		    cmp_one,
+	    cmp_exact}
+};
diff -urpN john-1.7.7.orig//src/cuda/cryptsha256.cu john-1.7.7-crypt-sha256_8//src/cuda/cryptsha256.cu
--- john-1.7.7.orig//src/cuda/cryptsha256.cu	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.7-crypt-sha256_8//src/cuda/cryptsha256.cu	2011-07-06 10:50:02.623207164 +0000
@@ -0,0 +1,331 @@
+/*
+* 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 <stdio.h>
+#include <stdlib.h>
+#include <assert.h>
+#include <string.h>
+#include "../cryptsha256.h"
+#define HANDLE_ERROR(err) (HandleError(err,__FILE__,__LINE__))
+
+static void HandleError(cudaError_t err, const char *file, int line)
+{
+	if (err != cudaSuccess) {
+		printf("%s in %s at line %d\n", cudaGetErrorString(err), file,
+		    line);
+		exit(EXIT_FAILURE);
+	}
+}
+
+extern "C" void sha256_crypt_gpu(crypt_sha256_password * inbuffer,
+    crypt_sha256_hash * outbuffer, crypt_sha256_salt * host_salt);
+
+__constant__ uint8_t dop[64] = { 0x80, 0 };
+
+__constant__ crypt_sha256_salt cuda_salt[1];
+__constant__ uint32_t 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
+};
+
+
+
+__device__ void init_ctx(sha256_ctx * ctx)
+{
+	ctx->H[0] = 0x6a09e667;
+	ctx->H[1] = 0xbb67ae85;
+	ctx->H[2] = 0x3c6ef372;
+	ctx->H[3] = 0xa54ff53a;
+	ctx->H[4] = 0x510e527f;
+	ctx->H[5] = 0x9b05688c;
+	ctx->H[6] = 0x1f83d9ab;
+	ctx->H[7] = 0x5be0cd19;
+	ctx->total = 0;
+	ctx->buflen = 0;
+}
+
+__device__ void insert_to_buffer(sha256_ctx * ctx, const uint8_t * string,
+    uint8_t len)
+{
+	int i = len;
+	uint8_t *d = &ctx->buffer[ctx->buflen];
+	while (i--)
+		*d++ = *string++;
+	ctx->buflen += len;
+}
+
+__device__ void sha256_block(sha256_ctx * ctx)
+{
+	int i;
+	uint32_t a = ctx->H[0];
+	uint32_t b = ctx->H[1];
+	uint32_t c = ctx->H[2];
+	uint32_t d = ctx->H[3];
+	uint32_t e = ctx->H[4];
+	uint32_t f = ctx->H[5];
+	uint32_t g = ctx->H[6];
+	uint32_t h = ctx->H[7];
+	uint32_t w[16];
+	uint32_t *data = (uint32_t *) ctx->buffer;
+
+#pragma unroll 16
+	for (i = 0; i < 16; i++)
+		w[i] = SWAP(data[i]);
+
+	uint32_t t1, t2;
+	for (i = 0; i < 16; i++) {
+		t1 = k[i] + w[i] + h + Sigma1(e) + Ch(e, f, g);
+		t2 = Maj(a, b, c) + Sigma0(a);
+
+		h = g;
+		g = f;
+		f = e;
+		e = d + t1;
+		d = c;
+		c = b;
+		b = a;
+		a = t1 + t2;
+	}
+
+	for (i = 16; i < 64; i++) {
+
+		w[i & 15] =
+		    sigma1(w[(i - 2) & 15]) + sigma0(w[(i - 15) & 15]) + w[(i -
+			16) & 15] + w[(i - 7) & 15];
+		t1 = k[i] + w[i & 15] + h + Sigma1(e) + Ch(e, f, g);
+		t2 = Maj(a, b, c) + Sigma0(a);
+
+		h = g;
+		g = f;
+		f = e;
+		e = d + t1;
+		d = c;
+		c = b;
+		b = a;
+		a = t1 + t2;
+	}
+
+	ctx->H[0] += a;
+	ctx->H[1] += b;
+	ctx->H[2] += c;
+	ctx->H[3] += d;
+	ctx->H[4] += e;
+	ctx->H[5] += f;
+	ctx->H[6] += g;
+	ctx->H[7] += h;
+
+}
+
+__device__ void ctx_update(sha256_ctx * ctx, const char *string, uint8_t len)
+{
+	ctx->total += len;
+	uint8_t startpos = ctx->buflen;
+	uint8_t partsize;
+	if (startpos + len <= 64) {
+		partsize = len;
+	} else
+		partsize = 64 - startpos;
+
+	insert_to_buffer(ctx, (const uint8_t *) string, partsize);
+	if (ctx->buflen == 64) {
+		uint8_t offset = 64 - startpos;
+		sha256_block(ctx);
+		ctx->buflen = 0;
+		insert_to_buffer(ctx, (const uint8_t *) (string + offset),
+		    len - offset);
+	}
+}
+
+/**
+  Add 0x80 byte to ctx->buffer and clean the rest of it
+**/
+__device__ void ctx_append_1(sha256_ctx * ctx)
+{
+	int i = 63 - ctx->buflen;
+	uint8_t *d = &ctx->buffer[ctx->buflen];
+	*d++ = 0x80;
+	while (i--)
+		*d++ = 0;
+}
+
+/**
+  Add ctx->bufflen at the end of ctx->buffer
+**/
+__device__ void ctx_add_length(sha256_ctx * ctx)
+{
+	uint32_t *blocks = (uint32_t *) ctx->buffer;
+	blocks[15] = SWAP(ctx->total * 8);
+}
+
+__device__ void finish_ctx(sha256_ctx * ctx)
+{
+	ctx_append_1(ctx);
+	ctx_add_length(ctx);
+	ctx->buflen = 0;
+}
+
+__device__ void clear_ctx_buffer(sha256_ctx * ctx)
+{
+	uint32_t *w = (uint32_t *) ctx->buffer;
+#pragma unroll 16
+	for (int i = 0; i < 16; i++)
+		w[i] = 0;
+	ctx->buflen = 0;
+
+}
+
+__device__ void sha256_digest(sha256_ctx * ctx, uint32_t * result)
+{
+	uint8_t i;
+	if (ctx->buflen <= 55) {	//data+0x80+datasize fits in one 512bit block
+		finish_ctx(ctx);
+		sha256_block(ctx);
+	} else {
+		uint8_t moved = 1;
+		if (ctx->buflen < 64) {	//data and 0x80 fits in one block
+			ctx_append_1(ctx);
+			moved = 0;
+		}
+		sha256_block(ctx);
+		clear_ctx_buffer(ctx);
+		if (moved)
+			ctx->buffer[0] = 0x80;	//append 1,the rest is already clean
+		ctx_add_length(ctx);
+		sha256_block(ctx);
+	}
+#pragma unroll 8
+	for (i = 0; i < 8; i++)
+		result[i] = SWAP(ctx->H[i]);
+}
+
+__device__ void sha256crypt(const char *pass, uint8_t passlength,
+    uint32_t * tresult, uint32_t idx, uint32_t rounds)
+{
+
+	uint32_t i, alt_result[8], temp_result[8];
+
+	sha256_ctx ctx, alt_ctx;
+	init_ctx(&ctx);
+	init_ctx(&alt_ctx);
+
+	ctx_update(&ctx, pass, passlength);
+	ctx_update(&ctx, cuda_salt[0].salt, cuda_salt[0].saltlen);
+
+	ctx_update(&alt_ctx, pass, passlength);
+	ctx_update(&alt_ctx, cuda_salt[0].salt, cuda_salt[0].saltlen);
+	ctx_update(&alt_ctx, pass, passlength);
+
+	sha256_digest(&alt_ctx, alt_result);
+
+	ctx_update(&ctx, (const char *) alt_result, passlength);
+
+
+	for (i = passlength; i > 0; i >>= 1) {
+		if ((i & 1) != 0)
+			ctx_update(&ctx, (const char *) alt_result, 32);
+		else
+			ctx_update(&ctx, pass, passlength);
+	}
+	sha256_digest(&ctx, alt_result);
+
+	init_ctx(&alt_ctx);
+	for (i = 0; i < passlength; i++)
+		ctx_update(&alt_ctx, pass, passlength);
+
+	sha256_digest(&alt_ctx, temp_result);
+
+	char p_sequence[32];
+	memcpy(p_sequence, temp_result, passlength);
+
+	init_ctx(&alt_ctx);
+	for (i = 0; i < 16 + ((unsigned char *) alt_result)[0]; i++)
+		ctx_update(&alt_ctx, cuda_salt[0].salt, cuda_salt[0].saltlen);
+
+	sha256_digest(&alt_ctx, temp_result);
+
+	uint8_t saltlength = cuda_salt[0].saltlen;
+
+	__shared__ char s_sequence[THREADS][16];
+	memcpy(s_sequence[threadIdx.x], temp_result, saltlength);
+
+	for (i = 0; i < rounds; i++) {
+		init_ctx(&ctx);
+
+		if ((i & 1) != 0)
+			ctx_update(&ctx, p_sequence, passlength);
+		else
+			ctx_update(&ctx, (const char *) alt_result, 32);
+
+		if ((i % 3) != 0)
+			ctx_update(&ctx, s_sequence[threadIdx.x], saltlength);
+
+		if ((i % 7) != 0)
+			ctx_update(&ctx, p_sequence, passlength);
+
+		if ((i & 1) != 0)
+			ctx_update(&ctx, (const char *) alt_result, 32);
+		else
+			ctx_update(&ctx, p_sequence, passlength);
+
+
+		sha256_digest(&ctx, alt_result);
+	}
+#pragma unroll 8
+	for (i = 0; i < 8; i++)
+		tresult[hash_addr(i, idx)] = alt_result[i];
+}
+
+__global__ void kernel_crypt_r(crypt_sha256_password * inbuffer,
+    uint32_t * outbuffer)
+{
+	uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+	sha256crypt((const char *) inbuffer[idx].v, inbuffer[idx].length,
+	    outbuffer, idx, cuda_salt[0].rounds);
+
+}
+
+
+void sha256_crypt_gpu(crypt_sha256_password * inbuffer,
+    crypt_sha256_hash * outbuffer, crypt_sha256_salt * host_salt)
+{
+
+	HANDLE_ERROR(cudaMemcpyToSymbol(cuda_salt, host_salt,
+		sizeof(crypt_sha256_salt)));
+
+	crypt_sha256_password *cuda_inbuffer;
+	uint32_t *cuda_outbuffer;
+	size_t insize = sizeof(crypt_sha256_password) * KEYS_PER_CRYPT;
+	size_t outsize = sizeof(crypt_sha256_hash) * KEYS_PER_CRYPT;
+	HANDLE_ERROR(cudaMalloc(&cuda_inbuffer, insize));
+	HANDLE_ERROR(cudaMalloc(&cuda_outbuffer, outsize));
+
+	HANDLE_ERROR(cudaMemcpy(cuda_inbuffer, inbuffer, insize,
+		cudaMemcpyHostToDevice));
+	dim3 dimGrid(BLOCKS);
+	dim3 dimBlock(THREADS);
+	kernel_crypt_r <<< dimGrid, dimBlock >>> (cuda_inbuffer,
+	    cuda_outbuffer);
+	cudaThreadSynchronize();
+	HANDLE_ERROR(cudaMemcpy(outbuffer, cuda_outbuffer, outsize,
+		cudaMemcpyDeviceToHost));
+
+	HANDLE_ERROR(cudaFree(cuda_inbuffer));
+	HANDLE_ERROR(cudaFree(cuda_outbuffer));
+}
diff -urpN john-1.7.7.orig//src/john.c john-1.7.7-crypt-sha256_8//src/john.c
--- john-1.7.7.orig//src/john.c	2011-02-27 12:31:36.000000000 +0000
+++ john-1.7.7-crypt-sha256_8//src/john.c	2011-06-21 10:53:25.000000000 +0000
@@ -39,6 +39,7 @@ extern int CPU_detect(void);
 
 extern struct fmt_main fmt_DES, fmt_BSDI, fmt_MD5, fmt_BF;
 extern struct fmt_main fmt_AFS, fmt_LM;
+extern struct fmt_main fmt_CRYPTSHA256CUDA;
 #ifdef HAVE_CRYPT
 extern struct fmt_main fmt_crypt;
 #endif
@@ -71,6 +72,7 @@ static void john_register_all(void)
 	john_register_one(&fmt_BF);
 	john_register_one(&fmt_AFS);
 	john_register_one(&fmt_LM);
+	john_register_one(&fmt_CRYPTSHA256CUDA);
 #ifdef HAVE_CRYPT
 	john_register_one(&fmt_crypt);
 #endif
diff -urpN john-1.7.7.orig//src/options.c john-1.7.7-crypt-sha256_8//src/options.c
--- john-1.7.7.orig//src/options.c	2011-02-25 08:40:09.000000000 +0000
+++ john-1.7.7-crypt-sha256_8//src/options.c	2011-06-21 10:53:25.000000000 +0000
@@ -105,7 +105,7 @@ static struct opt_entry opt_list[] = {
 "--salts=[-]COUNT           load salts with[out] at least COUNT passwords " \
 	"only\n" \
 "--format=NAME              force hash type NAME: " \
-	"DES/BSDI/MD5/BF/AFS/LM" MAYBE_CRYPT "\n" \
+	"DES/BSDI/MD5/BF/AFS/LM/CRYPTSHA256CUDA" MAYBE_CRYPT "\n" \
 "--save-memory=LEVEL        enable memory saving, at LEVEL 1..3\n"
 
 void opt_init(char *name, int argc, char **argv)
