diff -urpN magnumripper-magnum-jumbo-62b2557/doc/README-CUDA magnumripper-magnum-jumbo-62b2557-fast//doc/README-CUDA
--- magnumripper-magnum-jumbo-62b2557/doc/README-CUDA	2012-03-28 15:30:17.000000000 +0000
+++ magnumripper-magnum-jumbo-62b2557-fast//doc/README-CUDA	2012-03-31 17:24:57.651888344 +0000
@@ -10,5 +10,9 @@ Performance issues:
   If you have got Fermi or newer card change "-arch sm_10" to "-arch sm_20" in the NVCC_FLAGS (Makefile).
   Default THREADS and BLOCKS settings might not be optimal.
   To get better performance you can experiment with THREADS and BLOCKS macros defined for each format in cuda*.h file.
+  For MSCash2:
+  CARD NAME	BLOCKS	THREADS	SM	RESULT
+  GTX460	14	128	20	13398 c/s
+
 
 You can contact me at lukas[dot]odzioba[at]gmail[dot]com or john-dev mailing list
diff -urpN magnumripper-magnum-jumbo-62b2557/src/cuda/mscash2.cu magnumripper-magnum-jumbo-62b2557-fast//src/cuda/mscash2.cu
--- magnumripper-magnum-jumbo-62b2557/src/cuda/mscash2.cu	2012-03-28 15:30:17.000000000 +0000
+++ magnumripper-magnum-jumbo-62b2557-fast//src/cuda/mscash2.cu	2012-03-31 17:25:44.606988534 +0000
@@ -1,5 +1,5 @@
 /*
-* This software is Copyright (c) 2011 Lukas Odzioba <lukas dot odzioba at gmail dot com> 
+* This software is Copyright (c) 2011,2012 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.
 * Based on S3nf implementation http://openwall.info/wiki/john/MSCash2
@@ -247,8 +247,94 @@ __device__ void hmac_sha1(const uint8_t
 	PUT_WORD_32_BE(E, output, 16);
 }
 
+
+__device__ void big_hmac_sha1(
+    uint32_t * input, uint32_t inputlen, uint32_t * output,
+    uint32_t * ipad_state, uint32_t * opad_state,uint32_t *tmp_out)
+{
+	int i;
+	uint32_t temp, W[16]= {0};
+	uint32_t A, B, C, D, E;
+	uint8_t buf[64]= {0};
+	int lo;
+
+	for(i=0;i<5;i++)
+	  W[i]=SWAP(input[i]);
+	
+	uint32_t *src=(uint32_t*)buf;
+	i=64/4;
+	while(i--)
+	   *src++=0;
+	for(i=0;i<4;i++)
+	tmp_out[i]=SWAP(tmp_out[i]);
+	
+	for(lo=1; lo<ITERATIONS; lo++) {
+
+		A = ipad_state[0];
+		B = ipad_state[1];
+		C = ipad_state[2];
+		D = ipad_state[3];
+		E = ipad_state[4];
+
+		W[5]=0x80000000;
+		W[15]=0x2A0;
+		
+		#pragma unroll 9
+		for (i = 6; i < 15; i++) W[i]=0;
+
+		SHA1(A, B, C, D, E, W);
+
+		A += ipad_state[0];
+		B += ipad_state[1];
+		C += ipad_state[2];
+		D += ipad_state[3];
+		E += ipad_state[4];
+
+		W[0]=A;
+		W[1]=B;
+		W[2]=C;
+		W[3]=D;
+		W[4]=E;
+		W[5]=0x80000000;
+		W[15]=0x2A0;
+		#pragma unroll 9
+		for(i=6;i<15;i++) W[i]=0;
+			
+  
+		A = opad_state[0];
+		B = opad_state[1];
+		C = opad_state[2];
+		D = opad_state[3];
+		E = opad_state[4];
+
+		
+		SHA1(A, B, C, D, E, W);
+
+		A += opad_state[0];
+		B += opad_state[1];
+		C += opad_state[2];
+		D += opad_state[3];
+		E += opad_state[4];
+
+		W[0]=(A);
+		W[1]=(B);
+		W[2]=(C);
+		W[3]=(D);
+		W[4]=(E);
+	
+		tmp_out[0]^=(A);
+		tmp_out[1]^=(B);
+		tmp_out[2]^=(C);
+		tmp_out[3]^=(D);
+	}
+
+for(i=0;i<4;i++)
+	tmp_out[i]=SWAP(tmp_out[i]);
+}
+
+
 __device__ void pbkdf2(const uint8_t * pass, const uint8_t * salt,
-    int saltlen, uint8_t * out)
+            int saltlen, uint8_t * out)
 {
 	uint8_t temp[SHA1_DIGEST_LENGTH];
 	__shared__ uint8_t sbuf[THREADS][48];
@@ -256,12 +342,13 @@ __device__ void pbkdf2(const uint8_t * p
 	uint32_t ipad_state[5];
 	uint32_t opad_state[5];
 	int i, j;
-	uint8_t tmp_out[16];
+
+	uint32_t tmp_out[4];
 
 	i=48/4;
 	uint32_t *src=(uint32_t*)buf;
 	while(i--)
-	  *src++=0;
+		*src++=0;
 
 	memcpy(buf, salt, saltlen);
 	buf[saltlen + 3] = 0x01;
@@ -271,20 +358,19 @@ __device__ void pbkdf2(const uint8_t * p
 
 	hmac_sha1(pass, 16, buf, saltlen + 4, temp, ipad_state, opad_state);
 
+        uint32_t ftmp[5];
+   
 	memcpy(tmp_out, temp, 20);
+   	memcpy(ftmp, temp, 20);
 
-	for (i = 1; i < ITERATIONS; i++) {
-		hmac_sha1(pass, 16, temp, SHA1_DIGEST_LENGTH, temp, ipad_state,
-		    opad_state);
+	big_hmac_sha1( ftmp, SHA1_DIGEST_LENGTH, ftmp, ipad_state,
+	               opad_state,tmp_out);
 
-#pragma unroll 16
-		for (j = 0; j < 16; j++)
-			tmp_out[j] ^= temp[j];
-	}
-	memcpy(out, tmp_out, 20);
+	memcpy(out, tmp_out, 16);
 }
 
 
+
 __global__ void pbkdf2_kernel(mscash2_password * inbuffer,
     mscash2_hash * outbuffer)
 {