diff -urpN magnum-jumbo/src/cuda_phpass.h magnum-jumbo_phpassfixes//src/cuda_phpass.h
--- magnum-jumbo/src/cuda_phpass.h	2012-06-26 01:45:02.204701756 +0000
+++ magnum-jumbo_phpassfixes//src/cuda_phpass.h	2012-06-26 02:10:07.330452239 +0000
@@ -66,9 +66,6 @@
 #define AC4pCb				0xb18b7a77
 #define MASK1				0x77777777
 
-
-static char phpass_prefix[] = "$P$";
-
 typedef struct {
 	uint8_t v[15];
 	uint8_t length;
diff -urpN magnum-jumbo/src/cuda_phpass_fmt.c magnum-jumbo_phpassfixes//src/cuda_phpass_fmt.c
--- magnum-jumbo/src/cuda_phpass_fmt.c	2012-06-26 01:45:02.204701756 +0000
+++ magnum-jumbo_phpassfixes//src/cuda_phpass_fmt.c	2012-06-26 02:10:02.095451908 +0000
@@ -42,7 +42,7 @@ static struct fmt_tests tests[] = {
 	{"$P$900000000m6YEJzWtTmNBBL4jypbHv1", "openwall"},
 	{"$P$900000000zgzuX4Dc2091D8kak8RdR0", "h3ll00"},
 	{"$P$900000000qZTL5A0XQUX9hq0t8SoKE0", "1234567890"},
-	{"$P$900112200B9LMtPy2FSq910c1a6BrH0", "1234567890"},
+	{"$H$900112200B9LMtPy2FSq910c1a6BrH0", "1234567890"},
 	{"$P$900000000a94rg7R/nUK0icmALICKj1", "john"},
 	{"$P$9sadli2.wzQIuzsR2nYVhUSlHNKgG/0", "john"},
 	{"$P$90000000000tbNYOc9TwXvLEI62rPt1", ""},
@@ -89,10 +89,14 @@ static int valid(char *ciphertext, struc
 {
 	uint32_t i, count_log2;
 
+	int prefix=0;
 	if (strlen(ciphertext) != CIPHERTEXT_LENGTH)
 		return 0;
-	if (strncmp(ciphertext, phpass_prefix, 3) != 0)
-		return 0;
+	if (strncmp(ciphertext, "$P$", 3) == 0)
+		prefix=1;
+	if (strncmp(ciphertext, "$H$", 3) == 0)
+		prefix=1;
+	if(prefix==0) return 0;
 
 	for (i = 3; i < CIPHERTEXT_LENGTH; i++)
 		if (atoi64[ARCH_INDEX(ciphertext[i])] == 0x7F)
diff -urpN magnum-jumbo/src/opencl/phpass_kernel.cl magnum-jumbo_phpassfixes//src/opencl/phpass_kernel.cl
--- magnum-jumbo/src/opencl/phpass_kernel.cl	2012-06-26 01:45:02.214701640 +0000
+++ magnum-jumbo_phpassfixes//src/opencl/phpass_kernel.cl	2012-06-26 01:51:55.225702120 +0000
@@ -13,12 +13,14 @@ typedef struct {
 	unsigned int v[4];
 } phpass_hash;
 
-
-#define ROTATE_LEFT(x, s) ((x << s) | (x >> (32 - s)))
-#define F(x, y, z) (((x) & (y)) | ((~x) & (z)))
-#define G(x, y, z) (((x) & (z)) | ((y) & (~z)))
-#define H(x, y, z) ((x) ^ (y) ^ (z))
-#define I(x, y, z) ((y) ^ ((x) | (~z)))
+#define ROTATE_LEFT(x, s)	rotate(x,s)
+//#define F(x, y, z)		((z) ^ ((x) & ((y) ^ (z))))
+//#define G(x, y, z)		((y) ^ ((z) & ((x) ^ (y))))
+
+#define F(x, y, z) bitselect((z), (y), (x))
+#define G(x, y, z) bitselect((y), (x), (z))
+#define H(x, y, z)		((x) ^ (y) ^ (z))
+#define I(x, y, z)		((y) ^ ((x) | (~z)))
 
 
 #define FF(a, b, c, d, x, s, ac) \
@@ -69,7 +71,7 @@ typedef struct {
 
 
 
-inline void cuda_md5(char len,__private uint32_t * internal_ret,__private uint32_t * x)
+inline void md5(char len,__private uint32_t * internal_ret,__private uint32_t * x)
 {
 	x[len / 4] |= (((uint32_t) 0x80) << ((len & 0x3) << 3));
 	uint32_t x14 = len << 3;
@@ -187,7 +189,7 @@ __kernel void phpass
 		buff[i] = password[i - 8];
 	}
 
-	cuda_md5(8 + length, x, x);
+	md5(8 + length, x, x);
 	count = 1 << setting[SALT_SIZE+3];
 	for (i = 16; i < 16 + length; i++)
 		buff[i] = password[i - 16];
@@ -211,7 +213,6 @@ do {
 		c = 0x98badcfe;
 		d = 0x10325476;
 
-// FF(a, b, c, d, x0, S11, 0xd76aa478);
 		a = AC1 + x0;
 		a = ROTATE_LEFT(a, S11);
 		a += b;
diff -urpN magnum-jumbo/src/opencl_phpass_fmt.c magnum-jumbo_phpassfixes//src/opencl_phpass_fmt.c
--- magnum-jumbo/src/opencl_phpass_fmt.c	2012-06-26 01:45:02.216701656 +0000
+++ magnum-jumbo_phpassfixes//src/opencl_phpass_fmt.c	2012-06-26 02:10:40.624451722 +0000
@@ -168,11 +168,15 @@ static int valid(char *ciphertext, struc
 {
 	uint32_t i, j, count_log2, found;
 
+	int prefix=0;
 	if (strlen(ciphertext) != CIPHERTEXT_LENGTH)
 		return 0;
-	if (strncmp(ciphertext, phpass_prefix, 3) != 0)
-		return 0;
-
+	if (strncmp(ciphertext, "$P$", 3) == 0)
+		prefix=1;
+	if (strncmp(ciphertext, "$H$", 3) == 0)
+		prefix=1;
+	if(prefix==0) return 0;
+	
 	for (i = 3; i < CIPHERTEXT_LENGTH; i++) {
 		found = 0;
 		for (j = 0; j < 64; j++)