diff -urpN john-1.7.7.orig//README_CUDA john-1.7.7-crypt-sha256_3/README_CUDA
--- john-1.7.7.orig//README_CUDA	1970-01-01 00:00:00.000000000 +0000
+++ john-1.7.7-crypt-sha256_3/README_CUDA	2011-05-23 02:11:06.209368497 +0000
@@ -0,0 +1,25 @@
+README:
+This is 0-revision sha256 based crypt cuda patch for JtR 1.7.7
+
+This patch was tested on GeForce 9800GT card. 
+There is a big chance that you can improve performance on your machine only by different threads/blocks configuration.
+
+I had used 14 blocks, 192 threads and achieved:
+  Raw: 1518 c/s real, 1518 c/s virtual 
+
+You can try different settings by modification cryptsha256.h file:
+  #define BLOCKS 14
+  #define THREADS 192
+
+CRYPTSHA256CUDA 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
+
diff -urpN john-1.7.7.orig//src/Makefile john-1.7.7-crypt-sha256_3/src/Makefile
--- john-1.7.7.orig//src/Makefile	2011-02-27 12:31:50.000000000 +0000
+++ john-1.7.7-crypt-sha256_3/src/Makefile	2011-05-23 05:17:37.217361791 +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_3/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_3/src/cryptsha256.h	2011-05-23 05:13:19.024361718 +0000
@@ -0,0 +1,78 @@
+#ifndef _CRYPTSHA256_H
+#define _CRYPTSHA256_H
+#include <assert.h>
+#include "common.h"
+#include <stdint.h>
+#include <stdbool.h>
+
+#define uint32_t unsigned int
+
+#define BLOCKS 14
+#define THREADS 192
+
+#define KEYS_PER_CRYPT BLOCKS*THREADS
+
+#define P(x) { printf("%s\n","x"); fflush(stdout); }
+
+#define MAX(x,y) (x) > (y) ? (x) : (y)
+#define MIN(x,y) (x) < (y) ? (x) : (y)
+
+int validchar(char x)
+{
+	if (x == '.')
+		return 1;
+	if (x == '/')
+		return 1;
+	if (x >= '0' && x <= '9')
+		return 1;
+	if (x >= 'A' && x <= 'Z')
+		return 1;
+	if (x >= 'a' && x <= 'z')
+		return 1;
+	return 0;
+}
+
+typedef struct {
+	uint32_t rounds;
+	unsigned char saltoffset;
+	unsigned char saltlen;
+	char salt[58];
+} crypt_sha256_salt;
+
+typedef struct {
+	unsigned char length;
+	unsigned char v[63];	///more than we need but faster
+} crypt_sha256_password;
+
+typedef struct {
+	uint32_t v[8];		///128bit
+} crypt_sha256_hash;
+
+struct sha256_ctx {
+	uint32_t H[8];
+	uint32_t total[2];
+	uint32_t buflen;
+	char buffer[128];
+};
+
+#define Ch(x, y, z) ((x & y) ^ (~x & z))
+#define Maj(x, y, z) ((x & y) ^ (x & z) ^ (y & z))
+#define S0(x) (CYCLIC (x, 2) ^ CYCLIC (x, 13) ^ CYCLIC (x, 22))
+#define S1(x) (CYCLIC (x, 6) ^ CYCLIC (x, 11) ^ CYCLIC (x, 25))
+#define R0(x) (CYCLIC (x, 7) ^ CYCLIC (x, 18) ^ (x >> 3))
+#define R1(x) (CYCLIC (x, 17) ^ CYCLIC (x, 19) ^ (x >> 10))
+#define CYCLIC(w, s) ((w >> s) | (w << (32 - s)))
+
+#if __BYTE_ORDER == __LITTLE_ENDIAN
+# define SWAP(n) \
+    (((n) << 24) | (((n) & 0xff00) << 8) | (((n) >> 8) & 0xff00) | ((n) >> 24))
+#else
+# define SWAP(n) (n)
+#endif
+
+#define SALT_LEN_MAX 16
+#define ROUNDS_DEFAULT 5000
+#define ROUNDS_MIN 1000
+#define ROUNDS_MAX 999999999
+
+#endif
diff -urpN john-1.7.7.orig//src/cryptsha256cuda_fmt.c john-1.7.7-crypt-sha256_3/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_3/src/cryptsha256cuda_fmt.c	2011-05-23 05:17:08.245361950 +0000
@@ -0,0 +1,326 @@
+/*
+* This software is Copyright (c) 2011 Lukas Odzioba
+* <lukas dot odzioba at gmail dot com> 
+* Based on http://www.akkadia.org/drepper/SHA-crypt.txt
+* 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	54
+#define CIPHERTEXT_LENGTH	34	/// header = 3 | loopcnt = 1 | salt = 8 | ciphertext = 22
+#define BINARY_SIZE		32	///v
+#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,
+    crypt_sha256_hash * outbuffer, crypt_sha256_salt * host_salt);
+
+static crypt_sha256_password inbuffer[MAX_KEYS_PER_CRYPT];			/** plaintext ciphertexts **/
+static crypt_sha256_hash outbuffer[MAX_KEYS_PER_CRYPT];			/** calculated hashes **/
+
+static char currentsalt[64];
+crypt_sha256_salt host_salt;
+
+static struct fmt_tests tests[] = {
+	{"$5$saltstring$5B8vYYiY.CVt1RlTTf8KbXBH3hsxY/GNooZaBBGWEc5",
+	    "Hello world!"},
+	{"$5$rounds=1000$roundstoolow$yfvwcWrQ8l/K0DAWyuPMDNHpIVlTQebY9l/gL972bIC", "the minimum number is still observed"},
+	{"$5$rounds=5000$toolongsaltstrin$Un/5jzAHMgOGZ5.mWJpuVolil07guHPvOW8mGRcvxa5", "This is just a test"},
+	{"$5$saltstring$b6dP8GUbSwC1anz/QPmW0Zn16Hk.c0Z1a3MkmQIDiH4", "ala"},
+	{"$5$saltstring$0Az3qME7zTXm78kfHrR2OtT8WOu2gd8bcVn/9Y.3l/7", "john"},
+
+
+	{NULL}
+};
+
+static void init(void)
+{
+}
+
+static int valid(char *ciphertext)
+{
+	uint32_t i;
+	int len = strlen(ciphertext);
+
+	if (strncmp(ciphertext, "$5$", 3) != 0)
+		return 0;
+	char *p = strrchr(ciphertext, '$');
+	for (i = p - ciphertext + 1; i < len; i++)
+		if (!validchar(ARCH_INDEX(ciphertext[i])))
+			return 0;
+	if (len - (p - ciphertext + 1) != 43)
+		return 0;
+	return 1;
+};
+
+int findb64(char c)
+{
+	static const char b64t[64] =
+	    "./0123456789ABCDEFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuvwxyz";
+	int i = 0;
+	for (i = 0; i < 64; i++)
+		if (c == b64t[i])
+			return i;
+	puts("Unknown character in findb64");
+	exit(1);
+}
+
+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;
+	}
+	host_salt.saltoffset = offset;
+	memcpy(host_salt.salt, currentsalt + offset, 32);
+	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;
+}
+
+void save_hash(char *ciphertext, int x)
+{
+	char *b = (char *) outbuffer[x].v;
+	memset(b, 0, BINARY_SIZE);
+	char *p = (char *) strrchr(ciphertext, '$') + 1;
+	magic(p, b);
+}
+
+static void crypt_all(int count)
+{
+	sha256_crypt_gpu(inbuffer, outbuffer, &host_salt);
+}
+
+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_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_3/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_3/src/cuda/cryptsha256.cu	2011-05-23 05:17:10.933361796 +0000
@@ -0,0 +1,475 @@
+#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);
+	}
+}
+
+__device__ crypt_sha256_password *cuda_data = NULL;
+__device__ crypt_sha256_hash *cuda_data_out = NULL;
+__device__ char *cuda_setting = NULL;
+__device__ char *cuda_itoa64 = NULL;
+__device__ __constant__ crypt_sha256_salt cuda_salt[1];
+
+
+crypt_sha256_password *host_data = NULL;
+crypt_sha256_hash *host_data_out = NULL;
+
+extern "C" void sha256_crypt_gpu(crypt_sha256_password * inbuffer,
+    crypt_sha256_hash * outbuffer, crypt_sha256_salt * host_salt);
+
+__constant__ unsigned char fillbuf1[128] = { 0x80, 0 };
+
+__constant__ char sha256_salt_prefix[] = "$5$";
+__constant__ char sha256_rounds_prefix[] = "rounds=";
+__constant__ char b64t[65] =
+    "./0123456789ABCDEFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuvwxyz";
+__constant__ uint32_t K[64] = {
+	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 sha256_init_ctx_cu(struct sha256_ctx *ctx, uint32_t idx)
+{
+
+	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->total[1] = 0;
+	ctx->buflen = 0;
+}
+
+__device__ void
+sha256_process_block_cu(const void *buffer, size_t len, struct sha256_ctx *ctx,
+    uint32_t idx)
+{
+
+	register unsigned int t;
+	const uint32_t *words = (const uint32_t *) buffer;
+	size_t nwords = len / sizeof(uint32_t);
+
+	register uint32_t a = ctx->H[0];
+	register uint32_t b = ctx->H[1];
+	register uint32_t c = ctx->H[2];
+	register uint32_t d = ctx->H[3];
+	register uint32_t e = ctx->H[4];
+	register uint32_t f = ctx->H[5];
+	register uint32_t g = ctx->H[6];
+	register uint32_t h = ctx->H[7];
+
+	/* First increment the byte count.  FIPS 180-2 specifies the possible
+	   length of the file up to 2^64 bits.  Here we only compute the
+	   number of bytes.  Do a double word increment.  */
+	ctx->total[0] += len;
+	if (ctx->total[0] < len)
+		++ctx->total[1];
+
+	/* Process all bytes in the buffer with 64 bytes in each round of
+	   the loop.  */
+	while (nwords > 0) {
+		register uint32_t W[64];
+		register uint32_t a_save = a;
+		register uint32_t b_save = b;
+		register uint32_t c_save = c;
+		register uint32_t d_save = d;
+		register uint32_t e_save = e;
+		register uint32_t f_save = f;
+		register uint32_t g_save = g;
+		register uint32_t h_save = h;
+
+
+		/* Compute the message schedule according to FIPS 180-2:6.2.2 step 2.  */
+#pragma unroll 16
+		for (t = 0; t < 16; ++t) {
+			W[t] = SWAP(*words);
+			++words;
+		}
+#pragma unroll 48
+		for (t = 16; t < 64; ++t)
+			W[t] =
+			    R1(W[t - 2]) + W[t - 7] + R0(W[t - 15]) + W[t -
+			    16];
+
+		uint32_t T1, T2;
+
+#define E(T1,T2,a,b,c,d,e,f,g,h) \
+h = g;g = f;f = e;e = d + T1;\
+d = c;c = b;b = a;a = T1 + T2;
+
+#define F(T1,T2,a,b,c,d,e,f,g,h,t)\
+			T1 = h + S1(e) + Ch(e, f, g) + K[t] + W[t];\
+			T2 = S0(a) + Maj(a, b, c);\
+			E(T1,T2,a,b,c,d,e,f,g,h)
+
+		for (t = 0; t < 64; t++) {
+			F(T1, T2, a, b, c, d, e, f, g, h, t)
+		}
+
+		a += a_save;
+		b += b_save;
+		c += c_save;
+		d += d_save;
+		e += e_save;
+		f += f_save;
+		g += g_save;
+		h += h_save;
+
+		/* Prepare for the next round.  */
+		nwords -= 16;
+	}
+
+	/* Put checksum in context given as argument.  */
+	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 *sha256_finish_ctx_cu(struct sha256_ctx *ctx, void *resbuf,
+    uint32_t idx)
+{
+
+	/* Take yet unprocessed bytes into account.  */
+	uint32_t bytes = ctx->buflen;
+	size_t pad;
+
+	/* Now count remaining bytes.  */
+	ctx->total[0] += bytes;
+	if (ctx->total[0] < bytes)
+		++ctx->total[1];
+
+	pad = bytes >= 56 ? 64 + 56 - bytes : 56 - bytes;
+	memcpy(&ctx->buffer[bytes], fillbuf1, pad);
+	int i;
+
+	/* Put the 64-bit file length in *bits* at the end of the buffer.  */
+	*(uint32_t *) & ctx->buffer[bytes + pad + 4] =
+	    SWAP(ctx->total[0] << 3);
+	*(uint32_t *) & ctx->buffer[bytes + pad] =
+	    SWAP((ctx->total[1] << 3) | (ctx->total[0] >> 29));
+
+	/* Process last bytes.  */
+	sha256_process_block_cu(ctx->buffer, bytes + pad + 8, ctx, idx);
+
+	/* Put result from CTX in first 32 bytes following RESBUF.  */
+
+	unsigned char *cresbuf = (unsigned char *) resbuf;
+#pragma unroll 8
+	for (i = 0; i < 8; ++i) {
+		cresbuf[4 * i + 0] = (ctx->H[i] >> 24);
+		cresbuf[4 * i + 1] = (ctx->H[i] >> 16);
+		cresbuf[4 * i + 2] = (ctx->H[i] >> 8);
+		cresbuf[4 * i + 3] = (ctx->H[i]) & 0xff;
+	}
+	return resbuf;
+}
+
+__device__ void
+sha256_process_bytes_cu(const void *buffer, size_t len, struct sha256_ctx *ctx,
+    uint32_t idx)
+{
+	/* When we already have some bits in our internal buffer concatenate
+	   both inputs first.  */
+	if (ctx->buflen != 0) {
+		size_t left_over = ctx->buflen;
+		size_t add = 128 - left_over > len ? len : 128 - left_over;
+
+		memcpy(&ctx->buffer[left_over], buffer, add);
+		ctx->buflen += add;
+
+		if (ctx->buflen > 64) {
+			sha256_process_block_cu(ctx->buffer, ctx->buflen & ~63,
+			    ctx, idx);
+
+			ctx->buflen &= 63;
+			memcpy(ctx->buffer,
+			    &ctx->buffer[(left_over + add) & ~63],
+			    ctx->buflen);
+		}
+
+		buffer = (const char *) buffer + add;
+		len -= add;
+	}
+
+	/* Process available complete blocks.  */
+	if (len >= 64) {
+
+		while (len > 64) {
+			sha256_process_block_cu(memcpy(ctx->buffer,
+				buffer, 64), 64, ctx, idx);
+			buffer = (const char *) buffer + 64;
+			len -= 64;
+		}
+	}
+
+	/* Move remaining bytes into internal buffer.  */
+	if (len > 0) {
+		size_t left_over = ctx->buflen;
+
+		memcpy(&ctx->buffer[left_over], buffer, len);
+		left_over += len;
+		if (left_over >= 64) {
+			sha256_process_block_cu(ctx->buffer, 64, ctx, idx);
+			left_over -= 64;
+			memcpy(ctx->buffer, &ctx->buffer[64], left_over);
+		}
+		ctx->buflen = left_over;
+	}
+}
+
+
+__device__ int cudastrncmp(char *a, char *b, size_t bytes)
+{
+	int ret = 0;
+	int i = 0;
+	for (i = 0; i < bytes; i++) {
+		if (a[i] != b[i]) {
+			ret = 1;
+			break;
+		}
+	}
+	return ret;
+}
+
+__device__ size_t cuda_strlen(char *a)
+{
+	size_t ret = 0;
+	while (*a++ != 0)
+		ret++;
+	return ret;
+}
+
+__device__ size_t cudastrcspn(char *a, char b)
+{
+	int i, found = 0, len = cuda_strlen(a);
+	for (i = 0; i < len; i++)
+		if (a[i] == b) {
+			found = 1;
+			break;
+		}
+	if (found == 1)
+		return i;
+	return len;
+}
+
+__device__ void *cuda_mempcpy(void *dest, const void *src, size_t n)
+{
+	int i;
+	char *cd = (char *) dest;
+	char *cs = (char *) src;
+	for (i = 0; i < n; i++)
+		cd[i] = cs[i];
+	return &cd[n];
+}
+
+
+
+__device__ static char *sha256_crypt_r_cu(const char *key, size_t key_len,
+    char *buffer, int buflen, uint32_t idx)
+{
+
+
+	unsigned char alt_result[32];
+	unsigned char temp_result[32];
+	struct sha256_ctx ctx;
+	struct sha256_ctx alt_ctx;
+	size_t salt_len;
+	size_t cnt;
+	char *cp;
+	char *p_bytes;
+	char *s_bytes;
+
+	size_t rounds = cuda_salt[0].rounds;
+
+	char *salt = cuda_salt[0].salt;
+
+	salt_len = cuda_salt[0].saltlen;
+
+	/* Prepare for the real work.  */
+	sha256_init_ctx_cu(&ctx, idx);
+
+	/* Add the key string.  */
+	sha256_process_bytes_cu(key, key_len, &ctx, idx);
+
+	/* The last part is the salt string.  This must be at most 16
+	   characters and it ends at the first `$' character (for
+	   compatibility with existing implementations).  */
+	sha256_process_bytes_cu(salt, salt_len, &ctx, idx);
+
+
+	/* Compute alternate SHA256 sum with input KEY, SALT, and KEY.  The
+	   final result will be added to the first context.  */
+	sha256_init_ctx_cu(&alt_ctx, idx);
+
+	/* Add key.  */
+	sha256_process_bytes_cu(key, key_len, &alt_ctx, idx);
+
+	/* Add salt.  */
+	sha256_process_bytes_cu(salt, salt_len, &alt_ctx, idx);
+
+	/* Add key again.  */
+	sha256_process_bytes_cu(key, key_len, &alt_ctx, idx);
+
+	/* Now get result of this (32 bytes) and add it to the other
+	   context.  */
+	sha256_finish_ctx_cu(&alt_ctx, alt_result, idx);
+
+	/* Add for any character in the key one byte of the alternate sum.  */
+	for (cnt = key_len; cnt > 32; cnt -= 32)
+		sha256_process_bytes_cu(alt_result, 32, &ctx, idx);
+	sha256_process_bytes_cu(alt_result, cnt, &ctx, idx);
+
+	/* Take the binary representation of the length of the key and for every
+	   1 add the alternate sum, for every 0 the key.  */
+	for (cnt = key_len; cnt > 0; cnt >>= 1)
+		if ((cnt & 1) != 0)
+			sha256_process_bytes_cu(alt_result, 32, &ctx, idx);
+		else
+			sha256_process_bytes_cu(key, key_len, &ctx, idx);
+
+	/* Create intermediate result.  */
+	sha256_finish_ctx_cu(&ctx, alt_result, idx);
+
+	/* Start computation of P byte sequence.  */
+	sha256_init_ctx_cu(&alt_ctx, idx);
+
+	/* For every character in the password add the entire password.  */
+	for (cnt = 0; cnt < key_len; ++cnt)
+		sha256_process_bytes_cu(key, key_len, &alt_ctx, idx);
+
+	/* Finish the digest.  */
+	sha256_finish_ctx_cu(&alt_ctx, temp_result, idx);
+
+	/* Create byte sequence P.  */
+	char buf0[128];
+	cp = p_bytes = buf0;
+	for (cnt = key_len; cnt >= 32; cnt -= 32)
+		cp = (char *) cuda_mempcpy(cp, temp_result, 32);
+	memcpy(cp, temp_result, cnt);
+
+	/* Start computation of S byte sequence.  */
+	sha256_init_ctx_cu(&alt_ctx, idx);
+
+	/* For every character in the password add the entire password.  */
+	for (cnt = 0; cnt < 16 + alt_result[0]; ++cnt)
+		sha256_process_bytes_cu(salt, salt_len, &alt_ctx, idx);
+
+	/* Finish the digest.  */
+	sha256_finish_ctx_cu(&alt_ctx, temp_result, idx);
+
+	/* Create byte sequence S.  */
+
+	char *buf1 = &buf0[64];
+	cp = s_bytes = buf1;
+	for (cnt = salt_len; cnt >= 32; cnt -= 32)
+		cp = (char *) cuda_mempcpy(cp, temp_result, 32);
+	memcpy(cp, temp_result, cnt);
+
+	/* Repeatedly run the collected hash value through SHA256 to burn
+	   CPU cycles.  */
+	for (cnt = 0; cnt < rounds; ++cnt) {
+		/* New context.  */
+		sha256_init_ctx_cu(&ctx, idx);
+
+		/* Add key or last result.  */
+		if ((cnt & 1) != 0)
+			sha256_process_bytes_cu(p_bytes, key_len, &ctx, idx);
+		else
+			sha256_process_bytes_cu(alt_result, 32, &ctx, idx);
+
+		/* Add salt for numbers not divisible by 3.  */
+		if (cnt % 3 != 0)
+			sha256_process_bytes_cu(s_bytes, salt_len, &ctx, idx);
+
+		/* Add key for numbers not divisible by 7.  */
+		if (cnt % 7 != 0)
+			sha256_process_bytes_cu(p_bytes, key_len, &ctx, idx);
+
+		/* Add key or last result.  */
+		if ((cnt & 1) != 0)
+			sha256_process_bytes_cu(alt_result, 32, &ctx, idx);
+		else
+			sha256_process_bytes_cu(p_bytes, key_len, &ctx, idx);
+
+		/* Create intermediate result.  */
+		sha256_finish_ctx_cu(&ctx, alt_result, idx);
+	}
+
+
+	int i;
+#pragma unroll 32
+	for (i = 0; i < 32; i++)
+		buffer[i] = alt_result[i];
+	return buffer;
+}
+
+
+
+__global__ void kernel_crypt_r(crypt_sha256_password * inbuffer,
+    crypt_sha256_hash * outbuffer)
+{
+	uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+	sha256_crypt_r_cu((char *) inbuffer[idx].v, inbuffer[idx].length,
+	    (char *) outbuffer[idx].v, 128, idx);
+}
+ 
+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;
+	crypt_sha256_hash *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/formats.c john-1.7.7-crypt-sha256_3/src/formats.c
--- john-1.7.7.orig//src/formats.c	2011-02-27 12:34:10.000000000 +0000
+++ john-1.7.7-crypt-sha256_3/src/formats.c	2011-05-23 05:03:07.119361874 +0000
@@ -59,16 +59,16 @@ char *fmt_self_test(struct fmt_main *for
 
 		binary = format->methods.binary(ciphertext);
 		salt = format->methods.salt(ciphertext);
-
+	
 		if ((unsigned int)format->methods.salt_hash(salt) >=
 		    SALT_HASH_SIZE)
 			return "salt_hash";
 
 		format->methods.set_salt(salt);
 		format->methods.set_key(current->plaintext, index);
-
+		
 		format->methods.crypt_all(index + 1);
-
+		
 		for (size = 0; size < PASSWORD_HASH_SIZES; size++)
 		if (format->methods.binary_hash[size] &&
 		    format->methods.get_hash[size](index) !=
@@ -76,7 +76,7 @@ char *fmt_self_test(struct fmt_main *for
 			sprintf(s_size, "get_hash[%d](%d)", size, index);
 			return s_size;
 		}
-
+		
 		if (!format->methods.cmp_all(binary, index + 1))
 			return "cmp_all";
 		if (!format->methods.cmp_one(binary, index))
diff -urpN john-1.7.7.orig//src/john.c john-1.7.7-crypt-sha256_3/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_3/src/john.c	2011-05-01 14:22:41.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_3/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_3/src/options.c	2011-05-01 14:22:50.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)
