diff -urpN john-1.7.8.orig//doc/CRYPTSHA256_OPENCL patch_opencl/john-1.7.8-cryptsha256-opencl_0//doc/CRYPTSHA256_OPENCL
--- john-1.7.8.orig//doc/CRYPTSHA256_OPENCL	1970-01-01 00:00:00.000000000 +0000
+++ patch_opencl/john-1.7.8-cryptsha256-opencl_0//doc/CRYPTSHA256_OPENCL	2011-08-01 20:46:05.423966793 +0000
@@ -0,0 +1,59 @@
+  ,-*
+ (_)   Release Notes
+
+==== Version        ====
+    CRYPTSHA256 OPENCL ver. 0.0 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:
+	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     ====
+    Not tested on 64bit builds and AMD cards.
+
+==== Environment    ====
+    CRYPTSHA256OPENCL is supported only on following builds:
+      linux-x86-sse2, linux-x86-64
+    If you have problems while building patched JtR:
+
+    On nVidia cards:
+      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-cryptsha256opencl-0.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=cryptsha256opencl
+
+    Incremental attack:
+    $ john-1.7.8/run/john -i -format=cryptsha256opencl FILE
+
+==== Benchmarks     ==== 
+      ver 0:
+      GeForce GTX 460 1024MB   		Raw:	7027 c/s real, 7125 c/s virtual
+
diff -urpN john-1.7.8.orig//src/Makefile patch_opencl/john-1.7.8-cryptsha256-opencl_0//src/Makefile
--- john-1.7.8.orig//src/Makefile	2011-05-04 18:52:31.000000000 +0000
+++ patch_opencl/john-1.7.8-cryptsha256-opencl_0//src/Makefile	2011-08-01 20:42:43.936966843 +0000
@@ -19,6 +19,7 @@ OMPFLAGS =
 #OMPFLAGS = -fopenmp
 # Sun Studio with OpenMP (set the OMP_NUM_THREADS env var at runtime)
 #OMPFLAGS = -xopenmp
+OCLHEADERPATH = /usr/local/cuda/include/
 CFLAGS = -c -Wall -O2 -fomit-frame-pointer $(OMPFLAGS)
 ASFLAGS = -c $(OMPFLAGS)
 LDFLAGS = -s $(OMPFLAGS)
@@ -41,6 +42,9 @@ JOHN_OBJS = \
 	unafs.o \
 	unique.o
 
+OPENCL_OBJS = \
+	cryptsha256opencl_fmt.o
+
 BENCH_DES_OBJS_ORIG = \
 	DES_fmt.o DES_std.o
 
@@ -151,9 +155,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) $(OPENCL_OBJS) c3_fmt.o x86-64.o" \
 		CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \
-		LDFLAGS="$(LDFLAGS) -lcrypt"
+		LDFLAGS="$(LDFLAGS) -lcrypt -lOpenCL"
 
 linux-x86-64-avx:
 	$(LN) x86-64.h arch.h
@@ -188,9 +192,9 @@ 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) $(OPENCL_OBJS) c3_fmt.o x86.o x86-sse.o" \
 		CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \
-		LDFLAGS="$(LDFLAGS) -lcrypt"
+		LDFLAGS="$(LDFLAGS) -lcrypt -lOpenCL"
 
 linux-x86-mmx:
 	$(LN) x86-mmx.h arch.h
@@ -811,6 +815,9 @@ generic.h:
 		"$(BENCH_MD5_OBJS_DEPEND)" \
 		"$(BENCH_BF_OBJS_DEPEND)"
 
+cryptsha256opencl_fmt.o: cryptsha256opencl_fmt.c
+	$(CC) -I$(OCLHEADERPATH) $(CFLAGS) cryptsha256opencl_fmt.c
+
 bench: $(BENCH_OBJS)
 	$(LD) $(BENCH_OBJS) $(LDFLAGS) -o bench
 
@@ -890,6 +897,7 @@ depend:
 
 clean:
 	$(RM) $(PROJ) $(PROJ_DOS) $(PROJ_WIN32)
+	$(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/cryptsha256.cl patch_opencl/john-1.7.8-cryptsha256-opencl_0//src/cryptsha256.cl
--- john-1.7.8.orig//src/cryptsha256.cl	1970-01-01 00:00:00.000000000 +0000
+++ patch_opencl/john-1.7.8-cryptsha256-opencl_0//src/cryptsha256.cl	2011-08-01 20:37:35.213966741 +0000
@@ -0,0 +1,287 @@
+#include "cryptsha256opencl_fmt.h"
+
+__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
+};
+
+
+void init_ctx(__private 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;
+}
+
+
+
+void sha256_block(__private 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];
+	__private uint32_t w[16];
+	__private uint32_t *data = ctx->buffer;
+
+	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;
+
+}
+
+void insert_to_buffer(__private sha256_ctx * ctx, __private uint8_t * string,
+    uint8_t len)
+{
+	int i = len;
+	uint8_t *d = &ctx->buffer[ctx->buflen];
+	for (i = 0; i < len; i++)
+		d[i] = string[i];
+	ctx->buflen += len;
+}
+
+
+void ctx_update(__private sha256_ctx * ctx, __private 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, string, partsize);
+	if (ctx->buflen == 64) {
+		uint8_t offset = 64 - startpos;
+		sha256_block(ctx);
+		ctx->buflen = 0;
+		insert_to_buffer(ctx, (string + offset), len - offset);
+	}
+}
+
+/**
+  Add 0x80 byte to ctx->buffer and clean the rest of it
+**/
+void ctx_append_1(__private sha256_ctx * ctx)
+{
+	int i = 63 - ctx->buflen;
+	__private uint8_t *d = &ctx->buffer[ctx->buflen];
+	d[0] = 0x80;
+	d++;
+
+	while (i--) {
+		d[i] = 0;
+	}
+
+}
+
+/**
+  Add ctx->bufflen at the end of ctx->buffer
+**/
+void ctx_add_length(__private sha256_ctx * ctx)
+{
+	__private uint32_t *blocks = ctx->buffer;
+	blocks[15] = SWAP(ctx->total * 8);
+}
+
+void finish_ctx(__private sha256_ctx * ctx)
+{
+	ctx_append_1(ctx);
+	ctx_add_length(ctx);
+	ctx->buflen = 0;
+}
+
+void clear_ctx_buffer(__private sha256_ctx * ctx)
+{
+	__private uint32_t *w = ctx->buffer;
+	for (int i = 0; i < 16; i++)
+		w[i] = 0;
+	ctx->buflen = 0;
+
+}
+
+void sha256_digest(__private sha256_ctx * ctx, __private 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);
+	}
+	for (i = 0; i < 8; i++)
+		result[i] = SWAP(ctx->H[i]);
+}
+
+__kernel void cryptsha256
+    (__global const crypt_sha256_password * inbuffer,
+    __global crypt_sha256_hash * outbuffer,
+    __global const crypt_sha256_salt * hsalt) {
+	uint32_t i, idx = get_global_id(0);
+
+	__private uint8_t pass[PLAINTEXT_LENGTH];
+	__private uint8_t salt[SALT_LEN_MAX];
+
+	uint32_t saltlength = hsalt->saltlen;
+	uint32_t rounds = hsalt->rounds;
+	uint8_t passlength = inbuffer[idx].length;
+
+	for (i = 0; i < saltlength; i++)
+		salt[i] = hsalt->salt[i];
+	for (i = 0; i < passlength; i++)
+		pass[i] = inbuffer[idx].v[i];
+
+
+	__private uint32_t alt_result[8];
+	__private uint32_t temp_result[8];
+
+	__private sha256_ctx ctx, alt_ctx;
+	init_ctx(&ctx);
+	init_ctx(&alt_ctx);
+
+	ctx_update(&ctx, pass, passlength);
+	ctx_update(&ctx, salt, saltlength);
+
+	ctx_update(&alt_ctx, pass, passlength);
+	ctx_update(&alt_ctx, salt, saltlength);
+	ctx_update(&alt_ctx, pass, passlength);
+
+	sha256_digest(&alt_ctx, alt_result);
+
+	ctx_update(&ctx, alt_result, passlength);
+
+	for (i = passlength; i > 0; i >>= 1) {
+		if ((i & 1) != 0)
+			ctx_update(&ctx, 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);
+
+	__private char p_sequence[16];
+	__private char *src = temp_result;
+	for (i = 0; i < passlength; i++)
+		p_sequence[i] = src[i];
+
+	init_ctx(&alt_ctx);
+	for (i = 0; i < 16 + ((unsigned char *) alt_result)[0]; i++)
+		ctx_update(&alt_ctx, salt, saltlength);
+
+	sha256_digest(&alt_ctx, temp_result);
+
+	__private char s_sequence[16];
+	for (i = 0; i < saltlength; i++)
+		s_sequence[i] = src[i];
+
+	for (i = 0; i < rounds; i++) {
+		init_ctx(&ctx);
+
+		if ((i & 1) != 0)
+			ctx_update(&ctx, p_sequence, passlength);
+		else
+			ctx_update(&ctx, alt_result, 32);
+
+		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, alt_result, 32);
+		else
+			ctx_update(&ctx, p_sequence, passlength);
+
+
+		sha256_digest(&ctx, alt_result);
+	}
+
+
+	for (i = 0; i < 8; i++)
+		outbuffer[idx].v[i] = alt_result[i];
+}
diff -urpN john-1.7.8.orig//src/cryptsha256opencl_fmt.c patch_opencl/john-1.7.8-cryptsha256-opencl_0//src/cryptsha256opencl_fmt.c
--- john-1.7.8.orig//src/cryptsha256opencl_fmt.c	1970-01-01 00:00:00.000000000 +0000
+++ patch_opencl/john-1.7.8-cryptsha256-opencl_0//src/cryptsha256opencl_fmt.c	2011-08-01 20:35:39.873093379 +0000
@@ -0,0 +1,448 @@
+/*
+* 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 <CL/cl.h>
+
+#include "cryptsha256opencl_fmt.h"
+
+#define FORMAT_LABEL		"cryptsha256opencl"
+#define FORMAT_NAME		"CRYPTSHA256OPENCL"
+
+#define CRYPT_TYPE		"SHA256-based CRYPT"
+
+#define BENCHMARK_COMMENT	""
+#define BENCHMARK_LENGTH	-1
+
+#define BINARY_SIZE		32
+
+#define MIN_KEYS_PER_CRYPT	KEYS_PER_CRYPT
+#define MAX_KEYS_PER_CRYPT	KEYS_PER_CRYPT
+
+static crypt_sha256_password inbuffer[MAX_KEYS_PER_CRYPT];			/** plaintext ciphertexts **/
+static crypt_sha256_hash outbuffer[MAX_KEYS_PER_CRYPT];				/** calculated hashes **/
+
+static char source[SOURCE_BUFSIZE];
+static crypt_sha256_salt host_salt;
+
+static cl_context context;
+static cl_command_queue cq;
+static cl_program prog;
+static size_t src_size = 0;
+
+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"},
+
+	//{"$5$saltstring$7cz4bTeQ7MnNssphNhFVrITtuJYY/1tdvLL2uzLvOk8","a"},
+	//{"$5$saltstring$4Wjlxdm/Hbpo8ZQzKFazuvfUZPVVUQn6v1oPTX3nwX/","ab"},
+	//{"$5$saltstring$tDHA0KPsYQ8V.LDB1/fgW7cvROod5ZajSrx1tZU2JG9","abc"},
+	//{"$5$saltstring$LfhGTHVGfbAkxy/xKLgvSfXyeE7hZheoMRKhjfvNF6.","abcd"},
+	//{"$5$saltstring$Qg0Xm9f2VY.ePLAwNXnOPU/s8btLptK/tEU/gFnn8BD","abcde"},
+	//{"$5$saltstring$2Snf.yaHnLnLI3Qhsk2S119X4vKbwQyiTMOHp3Oy7F5","abcdef"},
+	//{"$5$saltstring$4Y5UR.6zwplRx6y93NJVyNkxqdlyT64EV68F2mCrZ16","abcdefg"},
+	{"$5$saltstring$bEM3iuUR.CTgy8Wygh4zu.CAgmlwx3uxm3dGA34.Ij4",
+	    "abcdefgh"},
+	//{"$5$saltstring$1/OrKXZSFlaEE2DKMhKKE8qCld5X0Ez0vtz5TvO3U3D","abcdefghi"},
+	//{"$5$saltstring$1IbZU70/Wo9m1b40ha6Ao8d.v6Ja0.bAFg5/QFVzoX/","abcdefghij"},
+
+	//{"$5$saltstring$S4gCgloAzqAXE5sRz9DShPvaXrwt4vjDJ4fYgIMbLo1","abcdefghijk"},
+	//{"$5$saltstring$AFNSzsWaoMDvt7lk2bx0rPapzCz2zGahXDdFeoXrNE9","abcdefghijkl"},
+	//{"$5$saltstring$QfHc8JBd2DfyloVL0YLDa23Dc67N9mbdYqyRJQlFqZ5","abcdefghijklm"},
+	//{"$5$saltstring$XKHiS.SSJ545PvJJr2t.HyUpmPZDAIT8fVvzr/HGhd0","abcdefghijklmn"},
+	//{"$5$saltstring$VxW44bFDcvixlQoTE4E.k5c8v1w0fGMyZ4tn8nGcWn0","abcdefghijklmno"},
+	{NULL}
+};
+
+static void readSource(char *src, const char *path)
+{
+	FILE *f = fopen(path, "r");
+	if (f != NULL) {
+		src_size = fread(src, 1, SOURCE_BUFSIZE, f);
+		fclose(f);
+	}
+}
+
+static void HandleError(cl_int cl_error, const char *file, int line)
+{
+	if (cl_error != CL_SUCCESS) {
+		printf("Error (%d) in file (%s) at line (%d)\n", cl_error,
+		    file, line);
+		exit(EXIT_FAILURE);
+	}
+}
+
+#define HANDLE_ERROR(cl_error) (HandleError(cl_error,__FILE__,__LINE__))
+
+
+static void init(void)
+{
+	readSource(source, SOURCE_NAME);
+
+	cl_int cl_error;
+	cl_platform_id platform;
+	cl_device_id device;
+	cl_uint platforms, devices;
+
+	HANDLE_ERROR(clGetPlatformIDs(1, &platform, &platforms));
+	HANDLE_ERROR(clGetDeviceIDs
+	    (platform, CL_DEVICE_TYPE_GPU, 1, &device, &devices));
+
+	cl_context_properties properties[] = {
+		CL_CONTEXT_PLATFORM, (cl_context_properties) platform,
+		0
+	};
+
+	char log[SOURCE_BUFSIZE];
+	HANDLE_ERROR(clGetDeviceInfo
+	    (device, CL_DEVICE_NAME, sizeof(log), log, NULL));
+	printf("\nOpenCL Device: %s\n", log);
+
+	context =
+	    clCreateContext(properties, 1, &device, NULL, NULL, &cl_error);
+	cq = clCreateCommandQueue(context, device, 0, &cl_error);
+
+	const char *srcptr[] = { source };
+	prog = clCreateProgramWithSource(context,
+	    1, srcptr, &src_size, &cl_error);
+
+	if (cl_error != CL_SUCCESS) {
+		printf("Error while creating program [%d]\n", cl_error);
+		exit(1);
+	}
+	cl_error = clBuildProgram(prog, 0, NULL, "", NULL, NULL);
+	if (cl_error != CL_SUCCESS) {
+		printf("Error while building program [%d]\n", cl_error);
+		cl_build_status build_status;
+		HANDLE_ERROR(clGetProgramBuildInfo(prog, device,
+			CL_PROGRAM_BUILD_STATUS,
+			sizeof(cl_build_status), &build_status, NULL));
+
+		char *build_log;
+		size_t ret_val_size;
+		HANDLE_ERROR(clGetProgramBuildInfo
+		    (prog, device, CL_PROGRAM_BUILD_LOG, 0, NULL,
+			&ret_val_size));
+
+		build_log = malloc(ret_val_size + 1);
+		HANDLE_ERROR(clGetProgramBuildInfo
+		    (prog, device, CL_PROGRAM_BUILD_LOG, ret_val_size,
+			build_log, NULL));
+		build_log[ret_val_size] = '\0';
+		printf("BUILD LOG: \n %s", build_log);
+		free(build_log);
+		exit(1);
+	}
+}
+
+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;
+};
+
+static int findb64(char c)
+{
+	return atoi64[(int) c] != 0x7F ? atoi64[(int) c] : 0;
+}
+
+static 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[64];
+	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)
+{
+	static char currentsalt[64];
+	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)
+{
+	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 crypt_all(int count)
+{
+	size_t insize = sizeof(crypt_sha256_password) * KEYS_PER_CRYPT;
+	size_t outsize = sizeof(crypt_sha256_hash) * KEYS_PER_CRYPT;
+	size_t saltsize = sizeof(crypt_sha256_salt);
+
+	///Alocate memory on the GPU
+	cl_int cl_error;
+	cl_mem mem_in, mem_out, mem_salt;
+	mem_in =
+	    clCreateBuffer(context, CL_MEM_READ_ONLY, insize, NULL, &cl_error);
+	mem_salt =
+	    clCreateBuffer(context, CL_MEM_READ_ONLY, saltsize, NULL,
+	    &cl_error);
+	mem_out =
+	    clCreateBuffer(context, CL_MEM_WRITE_ONLY, outsize, NULL,
+	    &cl_error);
+
+	///Assign kernel parameters
+	cl_kernel kernel = clCreateKernel(prog, KERNEL_NAME, &cl_error);
+	clSetKernelArg(kernel, 0, sizeof(mem_in), &mem_in);
+	clSetKernelArg(kernel, 1, sizeof(mem_out), &mem_out);
+	clSetKernelArg(kernel, 2, sizeof(mem_salt), &mem_salt);
+
+	///Copy data to GPU memory
+	cl_error =
+	    clEnqueueWriteBuffer(cq, mem_in, CL_FALSE, 0, insize, inbuffer, 0,
+	    NULL, NULL);
+	cl_error =
+	    clEnqueueWriteBuffer(cq, mem_salt, CL_FALSE, 0, saltsize,
+	    &host_salt, 0, NULL, NULL);
+
+	///Run kernel
+	size_t worksize = KEYS_PER_CRYPT;
+	size_t localworksize = THREADS;
+	cl_error =
+	    clEnqueueNDRangeKernel(cq, kernel, 1, NULL, &worksize,
+	    &localworksize, 0, NULL, NULL);
+	cl_error =
+	    clEnqueueReadBuffer(cq, mem_out, CL_FALSE, 0, outsize, outbuffer,
+	    0, NULL, NULL);
+
+	///Await completion of all the above
+	cl_error = clFinish(cq);
+}
+
+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;
+	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_CRYPTSHA256OPENCL = {
+	{
+		    FORMAT_LABEL,
+		    FORMAT_NAME,
+		    CRYPT_TYPE,
+		    BENCHMARK_COMMENT,
+		    BENCHMARK_LENGTH,
+		    PLAINTEXT_LENGTH,
+		    BINARY_SIZE,
+		    SALT_LENGTH,
+		    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/cryptsha256opencl_fmt.h patch_opencl/john-1.7.8-cryptsha256-opencl_0//src/cryptsha256opencl_fmt.h
--- john-1.7.8.orig//src/cryptsha256opencl_fmt.h	1970-01-01 00:00:00.000000000 +0000
+++ patch_opencl/john-1.7.8-cryptsha256-opencl_0//src/cryptsha256opencl_fmt.h	2011-08-01 20:38:51.791063533 +0000
@@ -0,0 +1,68 @@
+/*
+* 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
+
+#define BLOCKS 14
+#define THREADS 320
+
+#define KEYS_PER_CRYPT BLOCKS*THREADS
+
+#define SALT_LENGTH		16
+#define PLAINTEXT_LENGTH	15
+
+#define ROUNDS_DEFAULT 		5000
+#define ROUNDS_MIN 		1000
+#define ROUNDS_MAX 		999999999
+
+#define SOURCE_BUFSIZE 		1024*64
+#define SOURCE_NAME		"cryptsha256.cl"
+#define KERNEL_NAME		"cryptsha256"
+
+#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))
+
+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[SALT_LENGTH];
+} crypt_sha256_salt;
+
+typedef struct {
+	unsigned char length;
+	unsigned char v[PLAINTEXT_LENGTH];
+} crypt_sha256_password;
+
+
+#define hash_addr(j,idx) (((j)*(KEYS_PER_CRYPT))+(idx))
+typedef struct {
+	uint32_t v[8];
+} crypt_sha256_hash;
+
+#endif
diff -urpN john-1.7.8.orig//src/john.c patch_opencl/john-1.7.8-cryptsha256-opencl_0//src/john.c
--- john-1.7.8.orig//src/john.c	2011-02-27 12:31:36.000000000 +0000
+++ patch_opencl/john-1.7.8-cryptsha256-opencl_0//src/john.c	2011-08-01 07:20:39.995189385 +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_CRYPTSHA256OPENCL;
 #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_CRYPTSHA256OPENCL);
 #ifdef HAVE_CRYPT
 	john_register_one(&fmt_crypt);
 #endif
diff -urpN john-1.7.8.orig//src/options.c patch_opencl/john-1.7.8-cryptsha256-opencl_0//src/options.c
--- john-1.7.8.orig//src/options.c	2011-06-22 13:01:40.000000000 +0000
+++ patch_opencl/john-1.7.8-cryptsha256-opencl_0//src/options.c	2011-08-01 07:20:01.906966989 +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/CRYPTSHA256OPENCL" MAYBE_CRYPT "\n" \
 "--save-memory=LEVEL        enable memory saving, at LEVEL 1..3\n"
 
 void opt_init(char *name, int argc, char **argv)
