diff -urpN john-1.7.8.orig//doc/CRYPTSHA512_CUDA patch_cuda/john-1.7.8-crypt-sha512cuda_5//doc/CRYPTSHA512_CUDA
--- john-1.7.8.orig//doc/CRYPTSHA512_CUDA	1970-01-01 00:00:00.000000000 +0000
+++ patch_cuda/john-1.7.8-crypt-sha512cuda_5//doc/CRYPTSHA512_CUDA	2011-08-01 08:37:16.708967077 +0000
@@ -0,0 +1,60 @@
+  ,-*
+ (_)   Release Notes
+
+==== Version        ====
+    CRYPTSHA512 CUDA ver. 0.1 patch for John The Ripper 1.7.8
+
+==== Release Date   ====
+    DD.MM.RRRR = 01.08.2011
+
+==== Authors        ====
+    Lukas Odzioba, lukas[dot]odzioba[at]gmail[dot]com
+
+==== Change Log     ====
+      ver 0:
+	0.1 - Fixed CIPHERTEXT_LENGTH bug
+	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    ====
+    CRYPTSHA512CUDA is supported/tested only on following builds:
+      linux-x86-sse2, 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:
+	CUDAPATH = /usr/local/cuda/lib
+	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.8
+    $ cd john-1.7.8
+    $ patch -p1 < ../john-1.7.8-cryptsha512cuda-0.1.diff
+    $ cd john-1.7.8/src
+    for 32-bit systems:
+    $ make clean && make linux-x86-sse2
+    for 64-bit systems:
+    $ make clean && make linux-x86-64
+
+==== Usage          ====
+    Benchmark:
+    $ john-1.7.8/run/john -test --format=cryptsha512cuda
+
+    Incremental attack:
+    $ john-1.7.8/run/john -i -format=cryptsha512cuda FILE
+
+==== Benchmarks     ==== 
+    ver 0:
+      GeForce GTX 460 1024MB   		Raw:	4794 c/s real, 4794 c/s virtual
+
diff -urpN john-1.7.8.orig//src/Makefile patch_cuda/john-1.7.8-crypt-sha512cuda_5//src/Makefile
--- john-1.7.8.orig//src/Makefile	2011-05-04 18:52:31.000000000 +0000
+++ patch_cuda/john-1.7.8-crypt-sha512cuda_5//src/Makefile	2011-08-01 08:38:32.855217359 +0000
@@ -3,6 +3,7 @@
 # Copyright (c) 1996-2011 by Solar Designer
 #
 
+NVCC = nvccs
 CC = gcc
 AS = $(CC)
 LD = $(CC)
@@ -19,6 +20,9 @@ OMPFLAGS =
 #OMPFLAGS = -fopenmp
 # Sun Studio with OpenMP (set the OMP_NUM_THREADS env var at runtime)
 #OMPFLAGS = -xopenmp
+NVCCFLAGS = -c -Xptxas -v
+CUDAPATH = /usr/local/cuda/lib
+CUDA64PATH = /usr/local/cuda/lib64
 CFLAGS = -c -Wall -O2 -fomit-frame-pointer $(OMPFLAGS)
 ASFLAGS = -c $(OMPFLAGS)
 LDFLAGS = -s $(OMPFLAGS)
@@ -41,6 +45,9 @@ JOHN_OBJS = \
 	unafs.o \
 	unique.o
 
+CUDA_OBJS = \
+	cryptsha512cuda_fmt.o cryptsha512.o
+
 BENCH_DES_OBJS_ORIG = \
 	DES_fmt.o DES_std.o
 
@@ -151,9 +158,9 @@ default:
 linux-x86-64:
 	$(LN) x86-64.h arch.h
 	$(MAKE) $(PROJ) \
-		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86-64.o" \
+		JOHN_OBJS="$(JOHN_OBJS) $(CUDA_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
@@ -188,23 +195,23 @@ linux-x86-64-32-mmx:
 linux-x86-sse2:
 	$(LN) x86-sse.h arch.h
 	$(MAKE) $(PROJ) \
-		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86.o x86-sse.o" \
+		JOHN_OBJS="$(JOHN_OBJS) $(CUDA_OBJS) c3_fmt.o x86.o x86-sse.o" \
 		CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \
-		LDFLAGS="$(LDFLAGS) -lcrypt"
+		LDFLAGS="$(LDFLAGS) -L$(CUDAPATH) -lcudart -lcrypt"
 
 linux-x86-mmx:
 	$(LN) x86-mmx.h arch.h
 	$(MAKE) $(PROJ) \
-		JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86.o x86-mmx.o" \
+		JOHN_OBJS="$(JOHN_OBJS) $(CUDA_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" \
+		JOHN_OBJS="$(JOHN_OBJS) $(CUDA_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
@@ -811,6 +818,13 @@ generic.h:
 		"$(BENCH_MD5_OBJS_DEPEND)" \
 		"$(BENCH_BF_OBJS_DEPEND)"
 
+cryptsha512.o:  cryptsha512.h cuda/cryptsha512.cu
+	$(shell  cd cuda; nvcc $(NVCCFLAGS) cryptsha512.cu)
+	cp cuda/cryptsha512.o cryptsha512.o
+
+cryptsha512_fmt.o: cryptsha512.o cryptsha512cuda_fmt.c cryptsha512.h
+	$(CC)  $(CFLAGS) cryptsha512cuda_fmt.c
+
 bench: $(BENCH_OBJS)
 	$(LD) $(BENCH_OBJS) $(LDFLAGS) -o bench
 
@@ -890,6 +904,8 @@ depend:
 
 clean:
 	$(RM) $(PROJ) $(PROJ_DOS) $(PROJ_WIN32)
+	$(RM) cuda/cryptsha512.o cuda/*.cu~
+	$(RM) *~
 	$(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.8.orig//src/cryptsha512.h patch_cuda/john-1.7.8-crypt-sha512cuda_5//src/cryptsha512.h
--- john-1.7.8.orig//src/cryptsha512.h	1970-01-01 00:00:00.000000000 +0000
+++ patch_cuda/john-1.7.8-crypt-sha512cuda_5//src/cryptsha512.h	2011-08-01 08:27:20.935092594 +0000
@@ -0,0 +1,79 @@
+/*
+* 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 _CRYPTSHA512_H
+#define _CRYPTSHA512_H
+#include <assert.h>
+#include "common.h"
+#include <stdint.h>
+#include <stdbool.h>
+
+#define uint8_t  unsigned char
+#define uint32_t unsigned int
+#define uint64_t unsigned long long int
+
+#define BLOCKS 14
+#define THREADS 512
+
+#define KEYS_PER_CRYPT BLOCKS*THREADS
+
+#define MAX(x,y) (x) > (y) ? (x) : (y)
+#define MIN(x,y) (x) < (y) ? (x) : (y)
+#define SALT_LEN_MAX 16
+#define ROUNDS_DEFAULT 5000
+#define ROUNDS_MIN 1000
+#define ROUNDS_MAX 999999999
+
+static const char sha512_salt_prefix[] = "$6$";
+static const char sha256_rounds_prefix[] = "rounds=";
+
+
+#define rol(x,n) ((x << n) | (x >> (64-n)))
+#define ror(x,n) ((x >> n) | (x << (64-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,28))  ^ (ror(x,34)) ^ (ror(x,39)))
+#define Sigma1(x) ((ror(x,14))  ^ (ror(x,18)) ^ (ror(x,41)))
+#define sigma0(x) ((ror(x,1))  ^ (ror(x,8)) ^(x>>7))
+#define sigma1(x) ((ror(x,19)) ^ (ror(x,61)) ^(x>>6))
+
+# define SWAP32(n) \
+    (((n) << 24) | (((n) & 0xff00) << 8) | (((n) >> 8) & 0xff00) | ((n) >> 24))
+
+# define SWAP64(n) \
+  (((n) << 56)					\
+   | (((n) & 0xff00) << 40)			\
+   | (((n) & 0xff0000) << 24)			\
+   | (((n) & 0xff000000) << 8)			\
+   | (((n) >> 8) & 0xff000000)			\
+   | (((n) >> 24) & 0xff0000)			\
+   | (((n) >> 40) & 0xff00)			\
+   | ((n) >> 56))
+
+
+typedef struct {
+	uint64_t H[8];
+	uint32_t total;
+	uint32_t buflen;
+	uint8_t buffer[128];	//1024bits
+} sha512_ctx;
+
+typedef struct {
+	uint32_t rounds;
+	uint8_t saltlen;
+	char salt[63];
+} crypt_sha512_salt;
+
+typedef struct {
+	uint8_t length;
+	uint8_t v[16];
+} crypt_sha512_password;
+
+typedef struct {
+	uint64_t v[8];		//512 bits
+} crypt_sha512_hash;
+
+#endif
diff -urpN john-1.7.8.orig//src/cryptsha512cuda_fmt.c patch_cuda/john-1.7.8-crypt-sha512cuda_5//src/cryptsha512cuda_fmt.c
--- john-1.7.8.orig//src/cryptsha512cuda_fmt.c	1970-01-01 00:00:00.000000000 +0000
+++ patch_cuda/john-1.7.8-crypt-sha512cuda_5//src/cryptsha512cuda_fmt.c	2011-08-01 08:35:00.498216898 +0000
@@ -0,0 +1,343 @@
+/*
+* 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 "cryptsha512.h"
+
+#define FORMAT_LABEL		"cryptsha512cuda"
+#define FORMAT_NAME		"CRYPTSHA512CUDA"
+
+#define PHPASS_TYPE		"SHA512-based CRYPT"
+
+#define BENCHMARK_COMMENT	""
+#define BENCHMARK_LENGTH	-1
+
+#define PLAINTEXT_LENGTH	15
+#define CIPHERTEXT_LENGTH	34
+#define BINARY_SIZE		(3+16+86)
+#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
+
+static crypt_sha512_password inbuffer[MAX_KEYS_PER_CRYPT];			/** plaintext ciphertexts **/
+static crypt_sha512_hash outbuffer[MAX_KEYS_PER_CRYPT];			/** calculated hashes **/
+
+void sha512_crypt_gpu(crypt_sha512_password * inbuffer,
+    crypt_sha512_hash * outbuffer, crypt_sha512_salt * host_salt);
+
+static char currentsalt[64];
+static crypt_sha512_salt _salt;
+
+static struct fmt_tests tests[] = {
+
+	{"$6$saltstring$svn8UoSVapNtMuq1ukKS4tPQd8iKwSMHWjl/O817G3uBnIFNjnQJuesI68u4OTLiBFdcbYEdFCoEOfaS35inz1","Hello world!"},
+	//{"$6$saltstring$fgNTR89zXnDUV97U5dkWayBBRaB0WIBnu6s4T7T8Tz1SbUyewwiHjho25yWVkph2p18CmUkqXh4aIyjPnxdgl0","john"},
+	//{"$6$saltstring$MO53nAXQUKXVLlsbiXyPgMsR6q10N7eF7sPvanwdXnEeCj5kE3eYaRvFv0wVW1UZ4SnNTzc1v4OCOq1ASDQZY0","a"},
+	//{"$6$saltstring$q.eQ9PCFPe/tOHJPT7lQwnVQ9znjTT89hsg1NWHCRCAMsbtpBLbg1FLq7xo1BaCM0y/z46pXv4CGESVWQlOk30","ab"},
+	//{"$6$saltstring$pClZZISU0lxEwKr1z81EuJdiMLwWncjShXap25hiDGVMnCvlF5zS3ysvBdVRZqPDCdSTj06rwjrLX3bOS1Cak/","abc"},
+	//{"$6$saltstring$FJJAXr3hydAPJXM311wrzFhzheQ6LJHrufrYl2kBMnRD2pUi6jdS.fSBJ2J1Qfhcz9tPnlJOzeL7aIYi/dytg.","abcd"},
+	//{"$6$saltstring$XDecvJ/rq8tgbE1Pfuu1cTiZlhnbF5OA/vyP6HRPpDengVqhB38vbZTK/BDfPP6XBgvMzE.q9rj6Ck5blj/FK.","abcde"},
+	//{"$6$saltstring$hYPEYaHik6xSMGV1lDWhF0EerSUyCsC150POu9ksaftUWKWwV8TuqSeSLZUkUhjGy7cn.max5qd5IPSICeklL1","abcdef"},
+	//{"$6$saltstring$YBQ5J5EMRuC6k7B2GTsNaXx8u/957XMB.slQmY/lOjKd1zTIQF.ulLmy8O0VnJJ3cV.1pjP.KCgEjjMpz4pnS1","abcdefg"},
+	//{"$6$saltstring$AQapizZGhnIjtXF8OCvbSxQJBuOKvpzf1solf9b76wXFX0VRkqids5AC4YSibbhMSX0z4463sq1uAd9LvKNuO/","abcdefgh"},
+	//{"$6$saltstring$xc66FVXO.Zvv5pS02B4bCmJh5FCBAZpqTK3NoFxTU9U5b6BokbHwmeqQfMqrrkB3j9CXhCzgvC/pvoGPM1xgM1","abcdefghi"},
+	//{"$6$saltstring$Xet3A8EEzzSuL9hZZ31SfDVPT87qz3a.xxcH7eU50aqARlmXywdlfJ.6Cp/TFG1RcguqwrfUbZBbFn1BQ93Kv.","abcdefghij"},
+
+	//{"$6$saltstring$MeML1shJ8psyh5R9YJUZNYNqKzYeBvIsITqc/VqJfUDs8xO5YoUhCn4Db7CXuarMDVkBzIUfYq1d8Tj/T1WBU0","abcdefghijk"},
+	//{"$6$saltstring$i/3NHph8ZV2klLuOc5yX5kOnJWj9zuWbKiaa/NNEkYpNyamdQS1c7n2XQS3.B2Cs/eVyKwHf62PnOayqLLTOZ.","abcdefghijkl"},
+	//{"$6$saltstring$l2IxCS4o2S/vud70F1S5Z7H1WE67QFIXCYqskySdLFjjorEJdAnAp1ZqdgfNuZj2orjmeVDTsTXHpZ1IoxSKd.","abcdefghijklm"},
+	//{"$6$saltstring$PFzjspQs/CDXWALauDTav3u5bHB3n21xWrfwjnjpFO5eM5vuP0qKwDCXmlyZ5svEgsIH1oiZiGlRqkcBP5PiB.","abcdefghijklmn"},
+	//{"$6$saltstring$rdREv5Pd9C9YGtg.zXEQMb6m0sPeq4b6zFW9oWY9w4ZltmjH3yzMLgl9iBuez9DFFUvF5nJH3Y2xidiq1dH9M.", "abcdefghijklmno"},
+
+	{NULL}
+};
+
+static void init(void)
+{
+}
+
+static int valid(char *ciphertext)	///v
+{
+	uint32_t i, j;
+	int len = strlen(ciphertext);
+
+	if (strncmp(ciphertext, "$6$", 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;
+		if (found == 0) {
+			puts("not found");
+			return 0;
+		}
+	}
+	if (len - (p - ciphertext + 1) != 86)
+		return 0;
+	return 1;
+};
+
+static int findb64(char c)
+{
+	int i = 0;
+	for (i = 0; i < 64; i++)
+		if (c == itoa64[i])
+			return i;
+	return 0;
+}
+
+static void magic(char *crypt, unsigned 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, 21, 42);
+	_24bit_from_b64(4, 22, 43, 1);
+	_24bit_from_b64(8, 44, 2, 23);
+	_24bit_from_b64(12, 3, 24, 45);
+	_24bit_from_b64(16, 25, 46, 4);
+	_24bit_from_b64(20, 47, 5, 26);
+	_24bit_from_b64(24, 6, 27, 48);
+	_24bit_from_b64(28, 28, 49, 7);
+	_24bit_from_b64(32, 50, 8, 29);
+	_24bit_from_b64(36, 9, 30, 51);
+	_24bit_from_b64(40, 31, 52, 10);
+	_24bit_from_b64(44, 53, 11, 32);
+	_24bit_from_b64(48, 12, 33, 54);
+	_24bit_from_b64(52, 34, 55, 13);
+	_24bit_from_b64(56, 56, 14, 35);
+	_24bit_from_b64(60, 15, 36, 57);
+	_24bit_from_b64(64, 37, 58, 16);
+	_24bit_from_b64(68, 59, 17, 38);
+	_24bit_from_b64(72, 18, 39, 60);
+	_24bit_from_b64(76, 40, 61, 19);
+	_24bit_from_b64(80, 62, 20, 41);
+
+
+	uint32_t w = findb64(crypt[85]) << 6 | findb64(crypt[84]) << 0;
+	alt[63] = (w & 0xff);
+}
+
+
+static void *binary(char *ciphertext)
+{
+	static unsigned 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 (((uint64_t *) 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;
+	_salt.rounds = ROUNDS_DEFAULT;
+
+	if (strncmp((char *) "$6$", (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;
+			_salt.rounds =
+			    MAX(ROUNDS_MIN, MIN(srounds, ROUNDS_MAX));
+		}
+		offset = endp - currentsalt;
+	}
+	memcpy(_salt.salt, currentsalt + offset, 16);
+	_salt.saltlen = strlen(_salt.salt);
+}
+
+static void set_key(char *key, int index)
+{
+	int len = strlen(key);
+	inbuffer[index].length = len;
+	memcpy(inbuffer[index].v,key,len);
+}
+
+static char *get_key(int index)
+{
+	return (char *) inbuffer[index].v;
+}
+
+
+static void gpu_crypt_all(int count)
+{
+	sha512_crypt_gpu(inbuffer, outbuffer, &_salt);
+}
+
+static void crypt_all(int count)
+{
+	gpu_crypt_all(count);
+}
+
+static int get_hash_0(int index)
+{
+	return outbuffer[index].v[0] & 0xf;
+}
+
+static int get_hash_1(int index)
+{
+	return outbuffer[index].v[0] & 0xff;
+}
+
+static int get_hash_2(int index)
+{
+	return outbuffer[index].v[0] & 0xfff;
+}
+
+static int get_hash_3(int index)
+{
+	return outbuffer[index].v[0] & 0xffff;
+}
+
+static int get_hash_4(int index)
+{
+	return outbuffer[index].v[0] & 0xfffff;
+}
+
+static int cmp_all(void *binary, int count)
+{
+	uint32_t i;
+	uint64_t b = ((uint64_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;
+	uint64_t *t = (uint64_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_CRYPTSHA512CUDA = {
+	{
+		    FORMAT_LABEL,
+		    FORMAT_NAME,
+		    PHPASS_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,
+		    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.8.orig//src/cuda/cryptsha512.cu patch_cuda/john-1.7.8-crypt-sha512cuda_5//src/cuda/cryptsha512.cu
--- john-1.7.8.orig//src/cuda/cryptsha512.cu	1970-01-01 00:00:00.000000000 +0000
+++ patch_cuda/john-1.7.8-crypt-sha512cuda_5//src/cuda/cryptsha512.cu	2011-08-01 08:27:49.125341774 +0000
@@ -0,0 +1,359 @@
+/*
+* 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 "../cryptsha512.h"
+
+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);
+	}
+}
+
+#define HANDLE_ERROR(err) (HandleError(err,__FILE__,__LINE__))
+
+__constant__ uint64_t k[] = {
+	0x428a2f98d728ae22LL, 0x7137449123ef65cdLL, 0xb5c0fbcfec4d3b2fLL,
+	    0xe9b5dba58189dbbcLL,
+	0x3956c25bf348b538LL, 0x59f111f1b605d019LL, 0x923f82a4af194f9bLL,
+	    0xab1c5ed5da6d8118LL,
+	0xd807aa98a3030242LL, 0x12835b0145706fbeLL, 0x243185be4ee4b28cLL,
+	    0x550c7dc3d5ffb4e2LL,
+	0x72be5d74f27b896fLL, 0x80deb1fe3b1696b1LL, 0x9bdc06a725c71235LL,
+	    0xc19bf174cf692694LL,
+	0xe49b69c19ef14ad2LL, 0xefbe4786384f25e3LL, 0x0fc19dc68b8cd5b5LL,
+	    0x240ca1cc77ac9c65LL,
+	0x2de92c6f592b0275LL, 0x4a7484aa6ea6e483LL, 0x5cb0a9dcbd41fbd4LL,
+	    0x76f988da831153b5LL,
+	0x983e5152ee66dfabLL, 0xa831c66d2db43210LL, 0xb00327c898fb213fLL,
+	    0xbf597fc7beef0ee4LL,
+	0xc6e00bf33da88fc2LL, 0xd5a79147930aa725LL, 0x06ca6351e003826fLL,
+	    0x142929670a0e6e70LL,
+	0x27b70a8546d22ffcLL, 0x2e1b21385c26c926LL, 0x4d2c6dfc5ac42aedLL,
+	    0x53380d139d95b3dfLL,
+	0x650a73548baf63deLL, 0x766a0abb3c77b2a8LL, 0x81c2c92e47edaee6LL,
+	    0x92722c851482353bLL,
+	0xa2bfe8a14cf10364LL, 0xa81a664bbc423001LL, 0xc24b8b70d0f89791LL,
+	    0xc76c51a30654be30LL,
+	0xd192e819d6ef5218LL, 0xd69906245565a910LL, 0xf40e35855771202aLL,
+	    0x106aa07032bbd1b8LL,
+	0x19a4c116b8d2d0c8LL, 0x1e376c085141ab53LL, 0x2748774cdf8eeb99LL,
+	    0x34b0bcb5e19b48a8LL,
+	0x391c0cb3c5c95a63LL, 0x4ed8aa4ae3418acbLL, 0x5b9cca4f7763e373LL,
+	    0x682e6ff3d6b2b8a3LL,
+	0x748f82ee5defb2fcLL, 0x78a5636f43172f60LL, 0x84c87814a1f0ab72LL,
+	    0x8cc702081a6439ecLL,
+	0x90befffa23631e28LL, 0xa4506cebde82bde9LL, 0xbef9a3f7b2c67915LL,
+	    0xc67178f2e372532bLL,
+	0xca273eceea26619cLL, 0xd186b8c721c0c207LL, 0xeada7dd6cde0eb1eLL,
+	    0xf57d4f7fee6ed178LL,
+	0x06f067aa72176fbaLL, 0x0a637dc5a2c898a6LL, 0x113f9804bef90daeLL,
+	    0x1b710b35131c471bLL,
+	0x28db77f523047d84LL, 0x32caab7b40c72493LL, 0x3c9ebe0a15c9bebcLL,
+	    0x431d67c49c100d4cLL,
+	0x4cc5d4becb3e42b6LL, 0x597f299cfc657e2aLL, 0x5fcb6fab3ad6faecLL,
+	    0x6c44198c4a475817LL,
+};
+
+__constant__ crypt_sha512_salt cuda_salt[1];
+
+extern "C" void sha512_crypt_gpu(crypt_sha512_password * inbuffer,
+    crypt_sha512_hash * outbuffer, crypt_sha512_salt * salt);
+
+
+__device__ void init_ctx(sha512_ctx * ctx)
+{
+	ctx->H[0] = 0x6a09e667f3bcc908LL;
+	ctx->H[1] = 0xbb67ae8584caa73bLL;
+	ctx->H[2] = 0x3c6ef372fe94f82bLL;
+	ctx->H[3] = 0xa54ff53a5f1d36f1LL;
+	ctx->H[4] = 0x510e527fade682d1LL;
+	ctx->H[5] = 0x9b05688c2b3e6c1fLL;
+	ctx->H[6] = 0x1f83d9abfb41bd6bLL;
+	ctx->H[7] = 0x5be0cd19137e2179LL;
+	ctx->total = 0;
+	ctx->buflen = 0;
+}
+
+
+__device__ void insert_to_buffer(sha512_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 sha512_block(sha512_ctx * ctx)
+{
+	uint32_t i;
+	uint64_t a = ctx->H[0];
+	uint64_t b = ctx->H[1];
+	uint64_t c = ctx->H[2];
+	uint64_t d = ctx->H[3];
+	uint64_t e = ctx->H[4];
+	uint64_t f = ctx->H[5];
+	uint64_t g = ctx->H[6];
+	uint64_t h = ctx->H[7];
+
+
+	uint64_t w[16];
+
+	uint64_t *data = (uint64_t *) ctx->buffer;
+#pragma unroll 16
+	for (i = 0; i < 16; i++)
+		w[i] = SWAP64(data[i]);
+
+	uint64_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 < 80; 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_append_1(sha512_ctx * ctx)
+{
+	int i = 127 - ctx->buflen;
+	uint8_t *d = &ctx->buffer[ctx->buflen];
+	*d++ = 0x80;
+	while (i--)
+		*d++ = 0;
+}
+
+__device__ void ctx_add_length(sha512_ctx * ctx)
+{
+	uint64_t *blocks = (uint64_t *) ctx->buffer;
+	blocks[15] = SWAP64((uint64_t) ctx->total * 8);
+}
+
+__device__ void finish_ctx(sha512_ctx * ctx)
+{
+	ctx_append_1(ctx);
+	ctx_add_length(ctx);
+	ctx->buflen = 0;
+}
+
+
+__device__ void ctx_update(sha512_ctx * ctx, const char *string, uint8_t len)
+{
+	ctx->total += len;
+	uint8_t startpos = ctx->buflen;
+	uint8_t partsize;
+	if (startpos + len <= 128) {
+		partsize = len;
+	} else
+		partsize = 128 - startpos;
+
+	insert_to_buffer(ctx, (const uint8_t *) string, partsize);
+	if (ctx->buflen == 128) {
+		uint8_t offset = 128 - startpos;
+		sha512_block(ctx);
+		ctx->buflen = 0;
+		insert_to_buffer(ctx, (const uint8_t *) (string + offset),
+		    len - offset);
+	}
+}
+
+__device__ void clear_ctx_buffer(sha512_ctx * ctx)
+{
+	uint32_t *w = (uint32_t *) ctx->buffer;
+#pragma unroll 32
+	for (int i = 0; i < 32; i++)
+		w[i] = 0;
+	ctx->buflen = 0;
+
+}
+
+__device__ void sha512_digest(sha512_ctx * ctx, uint64_t * result)
+{
+	uint8_t i;
+	if (ctx->buflen <= 111) {	//data+0x80+datasize fits in one 1024bit block
+		finish_ctx(ctx);
+		sha512_block(ctx);
+	} else {
+		uint8_t moved = 1;
+		if (ctx->buflen < 128) {	//data and 0x80 fits in one block
+			ctx_append_1(ctx);
+			moved = 0;
+		}
+		sha512_block(ctx);
+		clear_ctx_buffer(ctx);
+		if (moved)
+			ctx->buffer[0] = 0x80;	//append 1,the rest is already clean
+		ctx_add_length(ctx);
+		sha512_block(ctx);
+	}
+#pragma unroll 8
+	for (i = 0; i < 8; i++)
+		result[i] = SWAP64(ctx->H[i]);
+}
+
+
+__device__ void sha512crypt(const char *pass, uint8_t passlength,
+    uint64_t * tresult, uint32_t idx, uint32_t rounds)
+{
+
+	uint64_t i, alt_result[8], temp_result[8];
+
+	sha512_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);
+
+	sha512_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, 64);
+		else
+			ctx_update(&ctx, pass, passlength);
+	}
+	sha512_digest(&ctx, alt_result);
+
+
+	init_ctx(&alt_ctx);
+	for (i = 0; i < passlength; i++)
+		ctx_update(&alt_ctx, pass, passlength);
+
+	sha512_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);
+
+	sha512_digest(&alt_ctx, temp_result);
+
+	uint8_t saltlength = cuda_salt[0].saltlen;
+
+	char s_sequence[16];
+	memcpy(s_sequence, 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, 64);
+
+		if ((i % 3) != 0)
+			ctx_update(&ctx, s_sequence, saltlength);
+
+		if ((i % 7) != 0)
+			ctx_update(&ctx, p_sequence, passlength);
+
+		if ((i & 1) != 0)
+			ctx_update(&ctx, (const char *) alt_result, 64);
+		else
+			ctx_update(&ctx, p_sequence, passlength);
+
+
+		sha512_digest(&ctx, alt_result);
+	}
+#pragma unroll 8
+	for (i = 0; i < 8; i++)
+		tresult[i] = alt_result[i];
+}
+
+
+
+__global__ void kernel_crypt_r(crypt_sha512_password * inbuffer,
+    crypt_sha512_hash * outbuffer)
+{
+	uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+	sha512crypt((const char *) inbuffer[idx].v, inbuffer[idx].length,
+	    outbuffer[idx].v, idx, cuda_salt[0].rounds);
+}
+
+void sha512_crypt_gpu(crypt_sha512_password * inbuffer,
+    crypt_sha512_hash * outbuffer, crypt_sha512_salt * host_salt)
+{
+
+	crypt_sha512_password *cuda_inbuffer;
+	crypt_sha512_hash *cuda_outbuffer;
+	size_t insize = sizeof(crypt_sha512_password) * KEYS_PER_CRYPT;
+	size_t outsize = sizeof(crypt_sha512_hash) * KEYS_PER_CRYPT;
+
+	HANDLE_ERROR(cudaMalloc(&cuda_inbuffer, insize));
+	HANDLE_ERROR(cudaMalloc(&cuda_outbuffer, outsize));
+
+	HANDLE_ERROR(cudaMemcpy(cuda_inbuffer, inbuffer, insize,
+		cudaMemcpyHostToDevice));
+	HANDLE_ERROR(cudaMemcpyToSymbol(cuda_salt, host_salt,
+		sizeof(crypt_sha512_salt)));
+
+	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.8.orig//src/john.c patch_cuda/john-1.7.8-crypt-sha512cuda_5//src/john.c
--- john-1.7.8.orig//src/john.c	2011-02-27 12:31:36.000000000 +0000
+++ patch_cuda/john-1.7.8-crypt-sha512cuda_5//src/john.c	2011-08-01 08:22:38.133966417 +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_CRYPTSHA512CUDA;
 #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_CRYPTSHA512CUDA);
 #ifdef HAVE_CRYPT
 	john_register_one(&fmt_crypt);
 #endif
diff -urpN john-1.7.8.orig//src/options.c patch_cuda/john-1.7.8-crypt-sha512cuda_5//src/options.c
--- john-1.7.8.orig//src/options.c	2011-06-22 13:01:40.000000000 +0000
+++ patch_cuda/john-1.7.8-crypt-sha512cuda_5//src/options.c	2011-08-01 08:22:05.613967010 +0000
@@ -100,7 +100,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/CRYPTSHA512CUDA" MAYBE_CRYPT "\n" \
 "--save-memory=LEVEL        enable memory saving, at LEVEL 1..3\n"
 
 void opt_init(char *name, int argc, char **argv)
