diff --git a/src/cuda/xsha512.cu b/src/cuda/xsha512.cu
index df3fcfa..011ceb3 100644
--- a/src/cuda/xsha512.cu
+++ b/src/cuda/xsha512.cu
@@ -141,6 +141,7 @@ __device__ void sha512_block(xsha512_ctx * ctx)
 	}
 
 	ctx->H[0] += a;
+#if 0
 	ctx->H[1] += b;
 	ctx->H[2] += c;
 	ctx->H[3] += d;
@@ -148,6 +149,7 @@ __device__ void sha512_block(xsha512_ctx * ctx)
 	ctx->H[5] += f;
 	ctx->H[6] += g;
 	ctx->H[7] += h;
+#endif
 }
 
 __device__ void xsha512_final(uint64_t *hash, xsha512_ctx *ctx)
@@ -182,10 +184,14 @@ __device__ void xsha512(const char* password, uint8_t pass_len, uint64_t *hash,
     xsha512_update(&ctx, password, pass_len);
     xsha512_final(hash, &ctx);
 
+#if 0
 	#pragma unroll 8
 	for(uint32_t i = 0; i < 8; ++i) {
 		hash[hash_addr(i, idx)] = SWAP64(ctx.H[i]);
 	}
+#else
+	hash[hash_addr(0, idx)] = SWAP64(ctx.H[0]);
+#endif
 }
 
 __global__ void kernel_xsha512(xsha512_key *cuda_password, xsha512_hash *cuda_hash)
diff --git a/src/cuda_xsha512.h b/src/cuda_xsha512.h
index 4e87ecc..3ee9c01 100644
--- a/src/cuda_xsha512.h
+++ b/src/cuda_xsha512.h
@@ -12,14 +12,22 @@
 #define uint32_t unsigned int
 #define uint64_t unsigned long long int
 
-#define BLOCKS 256
-#define THREADS 128
+#define BLOCKS 1024
+#define THREADS 480
 #define KEYS_PER_CRYPT BLOCKS*THREADS
 
 #define SALT_SIZE 4
+#if 0
 #define BINARY_SIZE 64
+#else
+#define BINARY_SIZE 8
+#endif
 
+#if 0
+#define PLAINTEXT_LENGTH		107
+#else
 #define PLAINTEXT_LENGTH		12
+#endif
 #define CIPHERTEXT_LENGTH		136
 
 extern uint8_t xsha512_key_changed;
@@ -64,7 +72,7 @@ typedef struct {
 } xsha512_key;
 
 typedef struct {
-    uint64_t v[8]; //512bits
+    uint64_t v[BINARY_SIZE / 8]; // up to 512 bits
 } xsha512_hash;
 
 #endif
diff --git a/src/cuda_xsha512_fmt.c b/src/cuda_xsha512_fmt.c
index f507689..e04dbef 100644
--- a/src/cuda_xsha512_fmt.c
+++ b/src/cuda_xsha512_fmt.c
@@ -1,6 +1,13 @@
 /*
- * This file is part of John the Ripper password cracker,
- * Copyright (c) 2008,2011 by Solar Designer
+ * Mac OS X 10.7+ salted SHA-512 password hashing, CUDA interface.
+ *
+ * Copyright (c) 2008,2011 Solar Designer (original CPU-only code)
+ * Copyright (c) 2012 myrice (interfacing to CUDA)
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted.
+ *
+ * There's ABSOLUTELY NO WARRANTY, express or implied.
  */
 
 
@@ -13,8 +20,8 @@
 #include "formats.h"
 
 #define FORMAT_LABEL			"xsha512-cuda"
-#define FORMAT_NAME			"Mac OS X 10.7+ salted SHA-512 CUDA"
-#define ALGORITHM_NAME			"64/" ARCH_BITS_STR
+#define FORMAT_NAME			"Mac OS X 10.7+ salted SHA-512"
+#define ALGORITHM_NAME			"CUDA"
 
 #define BENCHMARK_COMMENT		""
 #define BENCHMARK_LENGTH		0
@@ -220,9 +227,8 @@ static char *get_key(int index)
 
 static void crypt_all(int count)
 {
-    cuda_xsha512(gkey, &gsalt, ghash);
-	if(xsha512_key_changed)
-		xsha512_key_changed = 0;
+	cuda_xsha512(gkey, &gsalt, ghash);
+	xsha512_key_changed = 0;
 }
 
 static int cmp_all(void *binary, int count)
@@ -243,7 +249,7 @@ static int cmp_one(void *binary, int index)
 	int i;
 	uint64_t *b = (uint64_t *) binary;
 	uint64_t *t = (uint64_t *)ghash;
-	for (i = 0; i < 8; i++) {
+	for (i = 0; i < BINARY_SIZE / 8; i++) {
 		if (b[i] != t[hash_addr(i, index)])
 			return 0;
 	}
@@ -268,7 +274,7 @@ struct fmt_main fmt_cuda_xsha512 = {
 		SALT_SIZE,
 		MIN_KEYS_PER_CRYPT,
 		MAX_KEYS_PER_CRYPT,
-		FMT_CASE | FMT_8_BIT | FMT_OMP,
+		FMT_CASE | FMT_8_BIT,
 		tests
 	}, {
 		init,