diff -urpN magnum-jumbo/src/cuda/mscash.cu magnum-jumbo_mscash-prej7fix//src/cuda/mscash.cu --- magnum-jumbo/src/cuda/mscash.cu 2012-08-08 15:44:55.356433761 +0000 +++ magnum-jumbo_mscash-prej7fix//src/cuda/mscash.cu 2012-08-08 16:44:06.777807755 +0000 @@ -10,34 +10,9 @@ #include "../cuda_mscash.h" #include "cuda_common.cuh" -/* -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__)) -*/ -//extern "C" void mscash_init(int gpuid); extern "C" void cuda_mscash(mscash_password *, mscash_hash *, mscash_salt *); __constant__ mscash_salt cuda_salt[1]; -/*__host__ void mscash_init(int gpuid) -{ - int count; - HANDLE_ERROR(cudaGetDeviceCount(&count)); - if (gpuid < count) - cudaSetDevice(gpuid); - else { - printf("Invalid CUDA device id = %d\n", gpuid); - exit(1); - } -}*/ - __device__ static void md4_crypt(uint32_t * output, uint32_t * nt_buffer) { @@ -197,7 +172,7 @@ __global__ void mscash_kernel(mscash_pas mscash_hash * outbuffer) { uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x; - uint8_t *login = cuda_salt[0].salt; + uint8_t *login =(uint8_t*) cuda_salt[0].salt; uint8_t loginlength = cuda_salt[0].length; uint8_t *password = inbuffer[idx].v; diff -urpN magnum-jumbo/src/cuda/mscash2.cu magnum-jumbo_mscash-prej7fix//src/cuda/mscash2.cu --- magnum-jumbo/src/cuda/mscash2.cu 2012-08-08 15:44:55.356433761 +0000 +++ magnum-jumbo_mscash-prej7fix//src/cuda/mscash2.cu 2012-08-08 17:29:24.170583553 +0000 @@ -8,7 +8,8 @@ #include #include "../cuda_mscash2.h" #include "cuda_common.cuh" -extern "C" void mscash2_gpu(mscash2_password *, mscash2_hash *, mscash2_salt *); +extern "C" void mscash2_gpu(mscash2_password *, mscash2_hash *, + mscash2_salt *); __constant__ mscash2_salt cuda_salt[1]; @@ -150,7 +151,7 @@ __device__ __host__ void preproc(const u #pragma unroll 16 for (i = 0; i < 16; i++) GET_WORD_32_BE(W[i], ipad, i * 4); - + uint32_t A = INIT_A; uint32_t B = INIT_B; uint32_t C = INIT_C; @@ -174,12 +175,10 @@ __device__ void hmac_sha1(const uint8_t int i; uint32_t temp, W[16]; uint32_t A, B, C, D, E; - uint32_t state_A,state_B,state_C,state_D,state_E; + uint32_t state_A, state_B, state_C, state_D, state_E; uint8_t buf[64]; - uint32_t *src=(uint32_t*)buf; - i=64/4; - while(i--) - *src++=0; + for (i = 0; i < 64; i++) + buf[i] = 0; memcpy(buf, input, inputlen); buf[inputlen] = 0x80; @@ -190,12 +189,12 @@ __device__ void hmac_sha1(const uint8_t C = ipad_state[2]; D = ipad_state[3]; E = ipad_state[4]; - - state_A=A; - state_B=B; - state_C=C; - state_D=D; - state_E=E; + + state_A = A; + state_B = B; + state_C = C; + state_D = D; + state_E = E; for (i = 0; i < 16; i++) GET_WORD_32_BE(W[i], buf, i * 4); @@ -208,61 +207,50 @@ __device__ void hmac_sha1(const uint8_t D += state_D; E += state_E; - PUT_WORD_32_BE(A, buf, 0); - PUT_WORD_32_BE(B, buf, 4); - PUT_WORD_32_BE(C, buf, 8); - PUT_WORD_32_BE(D, buf, 12); - PUT_WORD_32_BE(E, buf, 16); - - buf[20] = 0x80; - PUT_WORD_32_BE(0x2A0, buf, 60); + W[0] = A; + W[1] = B; + W[2] = C; + W[3] = D; + W[4] = E; + W[5] = 0x80000000; + W[15] = 0x2A0; A = opad_state[0]; B = opad_state[1]; C = opad_state[2]; D = opad_state[3]; E = opad_state[4]; - - state_A=A; - state_B=B; - state_C=C; - state_D=D; - state_E=E; - - for (i = 0; i < 16; i++) - GET_WORD_32_BE(W[i], buf, i * 4); - SHA1(A, B, C, D, E, W); - - A += state_A; - B += state_B; - C += state_C; - D += state_D; - E += state_E; + SHA1_simply(A, B, C, D, E, W); - output[0]=SWAP(A); - output[1]=SWAP(B); - output[2]=SWAP(C); - output[3]=SWAP(D); - output[4]=SWAP(E); + A += opad_state[0]; + B += opad_state[1]; + C += opad_state[2]; + D += opad_state[3]; + E += opad_state[4]; + + output[0] = SWAP(A); + output[1] = SWAP(B); + output[2] = SWAP(C); + output[3] = SWAP(D); + output[4] = SWAP(E); } -__device__ void big_hmac_sha1( - uint32_t * input, uint32_t inputlen, - uint32_t * ipad_state, uint32_t * opad_state,uint32_t *tmp_out) +__device__ void big_hmac_sha1(uint32_t * input, uint32_t inputlen, + uint32_t * ipad_state, uint32_t * opad_state, uint32_t * tmp_out) { - int i,lo; + int i, lo; uint32_t temp, W[16]; uint32_t A, B, C, D, E; #pragma unroll 5 - for(i=0;i<5;i++) - W[i]=SWAP(input[i]); + for (i = 0; i < 5; i++) + W[i] = SWAP(input[i]); #pragma unroll 4 - for(i=0;i<4;i++) - tmp_out[i]=SWAP(tmp_out[i]); - - for(lo=1; losalt; uint32_t username_len = (uint32_t) host_salt->length; - //printf("username len=%d\n",username_len<<1); - int r=0; - if(username_len%2==1) - r=1; + int r = 0; + if (username_len % 2 == 1) + r = 1; for (i = 0; i < (username_len >> 1) + r; i++) ((uint32_t *) salt)[i] = username[2 * i] | (username[2 * i + 1] << 16); @@ -405,9 +374,7 @@ __host__ void mscash_cpu(mscash2_passwor buffer[i] = 0x80; buffer[14] = password_len << 4; - // printf("buffer[14]= %d\n",buffer[14]); md4_crypt(buffer, nt_hash); - //printf("buffer = %08x \n",((unsigned int *)buffer)[0]); memcpy((uint8_t *) nt_hash + 16, salt, username_len << 1); @@ -422,37 +389,28 @@ __host__ void mscash_cpu(mscash2_passwor nt_hash[14] = i << 4; md4_crypt(nt_hash, inbuffer[idx].dcc_hash); - } - } +} -__host__ void mscash2_gpu(mscash2_password * inbuffer, mscash2_hash * outbuffer, - mscash2_salt * host_salt) +__host__ void mscash2_gpu(mscash2_password * inbuffer, + mscash2_hash * outbuffer, mscash2_salt * host_salt) { - - mscash_cpu(inbuffer,outbuffer,host_salt); + + mscash_cpu(inbuffer, outbuffer, host_salt); mscash2_password *cuda_inbuffer; mscash2_hash *cuda_outbuffer; size_t insize = sizeof(mscash2_password) * KEYS_PER_CRYPT; size_t outsize = sizeof(mscash2_hash) * KEYS_PER_CRYPT; - + HANDLE_ERROR(cudaMemcpyToSymbol(cuda_salt, host_salt, sizeof(mscash2_salt))); - + HANDLE_ERROR(cudaMalloc(&cuda_inbuffer, insize)); HANDLE_ERROR(cudaMalloc(&cuda_outbuffer, outsize)); HANDLE_ERROR(cudaMemcpy(cuda_inbuffer, inbuffer, insize, cudaMemcpyHostToDevice)); - //int i; - //printf("usename len=%d dcc:\n",host_salt[0].length << 1); - //for(i=0;i<4;i++) - // printf("%08x ",inbuffer[0].dcc_hash[i]); - //puts(""); - //for(i=0;i<64;i++) - // printf("%d ",host_salt[0].unicode_salt[i]); - pbkdf2_kernel <<< BLOCKS, THREADS >>> (cuda_inbuffer, cuda_outbuffer); HANDLE_ERROR(cudaMemcpy(outbuffer, cuda_outbuffer, outsize, diff -urpN magnum-jumbo/src/cuda_mscash.h magnum-jumbo_mscash-prej7fix//src/cuda_mscash.h --- magnum-jumbo/src/cuda_mscash.h 2012-08-08 15:44:55.357683085 +0000 +++ magnum-jumbo_mscash-prej7fix//src/cuda_mscash.h 2012-08-08 17:30:45.952557917 +0000 @@ -51,7 +51,7 @@ typedef struct { typedef struct { uint8_t length; - uint8_t salt[SALT_LENGTH]; + char salt[SALT_LENGTH+1]; } mscash_salt; #endif diff -urpN magnum-jumbo/src/cuda_mscash2.h magnum-jumbo_mscash-prej7fix//src/cuda_mscash2.h --- magnum-jumbo/src/cuda_mscash2.h 2012-08-08 15:44:55.357683085 +0000 +++ magnum-jumbo_mscash-prej7fix//src/cuda_mscash2.h 2012-08-08 17:21:34.360683031 +0000 @@ -27,7 +27,7 @@ # define SWAP(n) \ (((n) << 24) | (((n) & 0xff00) << 8) | (((n) >> 8) & 0xff00) | ((n) >> 24)) -#define ITERATIONS 10240 +#define DEFAULT_ROUNDS 10240 #define INIT_A 0x67452301 #define INIT_B 0xefcdab89 @@ -79,24 +79,119 @@ ( W[t & 0x0F] = S(temp,1) ) \ ) -#define P1(a,b,c,d,e,x) \ +#define R16 \ +( \ + temp = W[2] ^ W[0], \ + ( W[0] = S(temp,1) ) \ +) + +#define R17 \ +( \ + temp = W[3] ^ W[1], \ + ( W[1] = S(temp,1) ) \ +) + +#define R18 \ +( \ + temp = W[15] ^ W[4] ^ W[2], \ + ( W[2] = S(temp,1) ) \ +) + +#define R19 \ +( \ + temp = W[0] ^ W[5] ^ W[3], \ + ( W[3] = S(temp,1) ) \ +) + +#define R20 \ +( \ + temp = W[1] ^ W[4], \ + ( W[4] = S(temp,1) ) \ +) + +#define R21 \ +( \ + temp = W[2] ^ W[5], \ + ( W[5] = S(temp,1) ) \ +) + +#define R22 \ +( \ + temp = W[3], \ + ( W[6] = S(temp,1) ) \ +) + +#define R23 \ +( \ + temp = W[4] ^ W[15], \ + ( W[7] = S(temp,1) ) \ +) + +#define R24 \ +( \ + temp = W[5] ^ W[0], \ + ( W[8] = S(temp,1) ) \ +) + +#define R25 \ +( \ + temp = W[6] ^ W[1], \ + ( W[9] = S(temp,1) ) \ +) + +#define R26 \ +( \ + temp = W[7] ^ W[2], \ + ( W[10] = S(temp,1) ) \ +) + +#define R27 \ +( \ + temp = W[8] ^ W[3], \ + ( W[11] = S(temp,1) ) \ +) + +#define R28 \ +( \ + temp = W[9] ^ W[4], \ + ( W[12] = S(temp,1) ) \ +) + +#define R29 \ +( \ + temp = W[10] ^ W[5] ^ W[15], \ + ( W[13] = S(temp,1) ) \ +) + +#define R30 \ +( \ + temp = W[11] ^ W[6] ^ W[0], \ + ( W[14] = S(temp,1) ) \ +) + +#define P1(a,b,c,d,e,x) \ +{ \ + e += S(a,5) + F1(b,c,d) + K1 + x; b = S(b,30); \ +} + +#define S1(a,b,c,d,e) \ { \ - e += S(a,5) + F1(b,c,d) + K1 + x; b = S(b,30); \ + e += S(a,5) + F1(b,c,d) + K1; b = S(b,30); \ } -#define P2(a,b,c,d,e,x) \ +#define P2(a,b,c,d,e,x) \ { \ - e += S(a,5) + F2(b,c,d) + K2 + x; b = S(b,30); \ + e += S(a,5) + F2(b,c,d) + K2 + x; b = S(b,30); \ } -#define P3(a,b,c,d,e,x) \ +#define P3(a,b,c,d,e,x) \ { \ - e += S(a,5) + F3(b,c,d) + K3 + x; b = S(b,30); \ + e += S(a,5) + F3(b,c,d) + K3 + x; b = S(b,30); \ } -#define P4(a,b,c,d,e,x) \ +#define P4(a,b,c,d,e,x) \ { \ - e += S(a,5) + F4(b,c,d) + K4 + x; b = S(b,30); \ + e += S(a,5) + F4(b,c,d) + K4 + x; b = S(b,30); \ } #define SHA1(A,B,C,D,E,W) \ @@ -181,6 +276,88 @@ P4(C, D, E, A, B, R(78));\ P4(B, C, D, E, A, R(79)); +#define SHA1_simply(A,B,C,D,E,W) \ + P1(A, B, C, D, E, W[0] );\ + P1(E, A, B, C, D, W[1] );\ + P1(D, E, A, B, C, W[2] );\ + P1(C, D, E, A, B, W[3] );\ + P1(B, C, D, E, A, W[4] );\ + P1(A, B, C, D, E, W[5] );\ + S1(E, A, B, C, D);\ + S1(D, E, A, B, C);\ + S1(C, D, E, A, B);\ + S1(B, C, D, E, A);\ + S1(A, B, C, D, E);\ + S1(E, A, B, C, D);\ + S1(D, E, A, B, C);\ + S1(C, D, E, A, B);\ + S1(B, C, D, E, A);\ + P1(A, B, C, D, E, W[15]);\ + P1(E, A, B, C, D, R16);\ + P1(D, E, A, B, C, R17);\ + P1(C, D, E, A, B, R18);\ + P1(B, C, D, E, A, R19);\ + P2(A, B, C, D, E, R20);\ + P2(E, A, B, C, D, R21);\ + P2(D, E, A, B, C, R22);\ + P2(C, D, E, A, B, R23);\ + P2(B, C, D, E, A, R24);\ + P2(A, B, C, D, E, R25);\ + P2(E, A, B, C, D, R26);\ + P2(D, E, A, B, C, R27);\ + P2(C, D, E, A, B, R28);\ + P2(B, C, D, E, A, R29);\ + P2(A, B, C, D, E, R30);\ + P2(E, A, B, C, D, R(31));\ + P2(D, E, A, B, C, R(32));\ + P2(C, D, E, A, B, R(33));\ + P2(B, C, D, E, A, R(34));\ + P2(A, B, C, D, E, R(35));\ + P2(E, A, B, C, D, R(36));\ + P2(D, E, A, B, C, R(37));\ + P2(C, D, E, A, B, R(38));\ + P2(B, C, D, E, A, R(39));\ + P3(A, B, C, D, E, R(40));\ + P3(E, A, B, C, D, R(41));\ + P3(D, E, A, B, C, R(42));\ + P3(C, D, E, A, B, R(43));\ + P3(B, C, D, E, A, R(44));\ + P3(A, B, C, D, E, R(45));\ + P3(E, A, B, C, D, R(46));\ + P3(D, E, A, B, C, R(47));\ + P3(C, D, E, A, B, R(48));\ + P3(B, C, D, E, A, R(49));\ + P3(A, B, C, D, E, R(50));\ + P3(E, A, B, C, D, R(51));\ + P3(D, E, A, B, C, R(52));\ + P3(C, D, E, A, B, R(53));\ + P3(B, C, D, E, A, R(54));\ + P3(A, B, C, D, E, R(55));\ + P3(E, A, B, C, D, R(56));\ + P3(D, E, A, B, C, R(57));\ + P3(C, D, E, A, B, R(58));\ + P3(B, C, D, E, A, R(59));\ + P4(A, B, C, D, E, R(60));\ + P4(E, A, B, C, D, R(61));\ + P4(D, E, A, B, C, R(62));\ + P4(C, D, E, A, B, R(63));\ + P4(B, C, D, E, A, R(64));\ + P4(A, B, C, D, E, R(65));\ + P4(E, A, B, C, D, R(66));\ + P4(D, E, A, B, C, R(67));\ + P4(C, D, E, A, B, R(68));\ + P4(B, C, D, E, A, R(69));\ + P4(A, B, C, D, E, R(70));\ + P4(E, A, B, C, D, R(71));\ + P4(D, E, A, B, C, R(72));\ + P4(C, D, E, A, B, R(73));\ + P4(B, C, D, E, A, R(74));\ + P4(A, B, C, D, E, R(75));\ + P4(E, A, B, C, D, R(76));\ + P4(D, E, A, B, C, R(77));\ + P4(C, D, E, A, B, R(78));\ + P4(B, C, D, E, A, R(79)); + static const char mscash2_prefix[] = "$DCC2$"; typedef struct { @@ -197,6 +374,7 @@ typedef struct { uint8_t length; uint8_t salt[19]; uint8_t unicode_salt[64]; + uint32_t rounds; } mscash2_salt; #endif diff -urpN magnum-jumbo/src/cuda_mscash2_fmt.c magnum-jumbo_mscash-prej7fix//src/cuda_mscash2_fmt.c --- magnum-jumbo/src/cuda_mscash2_fmt.c 2012-08-08 15:44:55.357683085 +0000 +++ magnum-jumbo_mscash-prej7fix//src/cuda_mscash2_fmt.c 2012-08-08 17:29:03.909808103 +0000 @@ -15,8 +15,8 @@ #define FORMAT_LABEL "mscash2-cuda" #define FORMAT_NAME "M$ Cache Hash 2 (DCC2) PBKDF2-HMAC-SHA-1" -#define MAX_CIPHERTEXT_LENGTH (7+19+32) -#define ALGORITHM_NAME "CUDA, unreliable, may miss guesses" +#define MAX_CIPHERTEXT_LENGTH (8+5+19+32) +#define ALGORITHM_NAME "CUDA" #define BENCHMARK_COMMENT "" #define BENCHMARK_LENGTH -1 @@ -27,37 +27,19 @@ static mscash2_hash *outbuffer; static mscash2_salt currentsalt; static struct fmt_tests tests[] = { - {"$DCC2$test#a86012faf7d88d1fc037a69764a92cac", "password"}, - {"$DCC2$test#a86012faf7d88d1fc037a69764a92cac", "password"}, - {"$DCC2$#59137848828d14b1fca295a5032b52a1", "a" }, //Empty Salt - {"$DCC2$administrator#a150f71752b5d605ef0b2a1e98945611","a"}, - {"$DCC2$administrator#c14eb8279e4233ec14e9d393637b65e2","ab"}, - {"$DCC2$administrator#8ce9c0279b4e6f226f52d559f9c2c5f3","abc"}, - {"$DCC2$administrator#2fc788d09fad7e26a92d12356fa44bdf","abcd"}, - {"$DCC2$administrator#6aa19842ffea11f0f0c89f8ca8d245bd","abcde"}, - {"$DCC2$test#a86012faf7d88d1fc037a69764a92cac", "password"}, - {"$DCC2$test3#360e51304a2d383ea33467ab0b639cc4", "test3" }, - {"$DCC2$test4#6f79ee93518306f071c47185998566ae", "test4" }, - {"$DCC2$january#26b5495b21f9ad58255d99b5e117abe2", "verylongpassword" }, - {"$DCC2$february#469375e08b5770b989aa2f0d371195ff", "(##)(&#*%%" }, - {"$DCC2$TEST2#c6758e5be7fc943d00b97972a8a97620", "test2" }, // salt is lowercased before hashing - {"$DCC2$john#ef9a549b7077f12143c18aecb8487d68","w00t"}, - - //{"$DCC2$administrator#56f8c24c5a914299db41f70e9b43f36d", "w00t" }, - //{"$DCC2$AdMiNiStRaToR#56f8C24c5A914299Db41F70e9b43f36d", "w00t" }, //Salt and hash are lowercased - - /*{"$DCC2$nineteen_characters#c4201b8267d74a2db1d5d19f5c9f7b57", "verylongpassword" }, //max salt_length - {"$DCC2$nineteen_characters#87136ae0a18b2dafe4a41d555425b2ed", "w00t"}, -*/// - //{"$DCC2$eighteencharacters#fc5df74eca97afd7cd5abb0032496223", "w00t" }, - //{"$DCC2$john-the-ripper#495c800a038d11e55fafc001eb689d1d", "batman#$@#1991" }, -/// {"$DCC2$jack#dc70386d419fc48442e6d0f64fa5f3da","Skipping and& Dipping"}, //passlen = 21 -/// {"$DCC2$john#d089ffa28f7508f67dcbb85f46b26886","0123456789012345678901234"}, //passlen =25 -/// {"$DCC2$john#51708f1b4587d6e0fb62f71b256692b5","012345678901234567890123456"}, //passlen 27 - //{"$DCC2$john#e7eb0fe73504d06f796615e6de083963","0123456789012345678901234567"}, //passlen 28 - //{"$DCC2$john#799c528e18017c5bc2e8d272de8e94ba","012345678901234567890123456789"}, //passlen= 30 - //{"$DCC2$john#caf1f2deef864f10a67c30be23087fa1","012345678901234567890123456789_"}, //passlen = 31 - + {"$DCC2$10240#test1#607bbe89611e37446e736f7856515bf8", "test1"}, + {"$DCC2$10240#Joe#e09b38f84ab0be586b730baf61781e30", "qerwt"}, + {"$DCC2$10240#Joe#6432f517a900b3fc34ffe57f0f346e16", "12345"}, + {"c0cbe0313a861062e29f92ede58f9b36", "", {"bin"}}, // nullstring password + {"87136ae0a18b2dafe4a41d555425b2ed", "w00t", {"nineteen_characters"}}, // max salt length + {"fc5df74eca97afd7cd5abb0032496223", "w00t", {"eighteencharacters"}}, +//unsupported salts lengths +// {"cfc6a1e33eb36c3d4f84e4c2606623d2", "longpassword", {"twentyXXX_characters"} }, +// {"99ff74cea552799da8769d30b2684bee", "longpassword", {"twentyoneX_characters"} }, +// {"0a721bdc92f27d7fb23b87a445ec562f", "longpassword", {"twentytwoXX_characters"} }, + {"$DCC2$10240#TEST2#c6758e5be7fc943d00b97972a8a97620", "test2"}, // salt is lowercased before hashing + {"$DCC2$10240#test3#360e51304a2d383ea33467ab0b639cc4", "test3"}, + {"$DCC2$10240#test4#6f79ee93518306f071c47185998566ae", "test4"}, {NULL} }; @@ -86,28 +68,38 @@ static void init(struct fmt_main *self) static int valid(char *ciphertext, struct fmt_main *self) { - char *pos,*hash = strrchr(ciphertext, '#') + 1; - int hashlength = 0; - int saltlength = 0; +#ifdef _MSCASH2_DEBUG + printf("valid(%s)\n", ciphertext); +#endif + int i, l = strlen(ciphertext), saltlength = 0; if (strncmp(ciphertext, mscash2_prefix, strlen(mscash2_prefix)) != 0) return 0; - - if (hash == NULL) + if (l <= 32 || l > MAX_CIPHERTEXT_LENGTH) return 0; - while (hash < ciphertext + strlen(ciphertext)) { - if (atoi16[ARCH_INDEX(*hash++)] == 0x7f) - return 0; - hashlength++; - } - if (hashlength != 32) + l -= 32; + if (ciphertext[l - 1] != '#') return 0; - - pos = ciphertext + strlen(mscash2_prefix); - while (*pos++ != '#') { - if (saltlength == 19) + for (i = l; i < l + 32; i++) + if (atoi16[ARCH_INDEX(ciphertext[i])] == 0x7F) return 0; - saltlength++; + + i = 6; + while (ciphertext[i] && ciphertext[i] != '#') + i++; + i++; + while (ciphertext[i] && ciphertext[i] != '#') + i++, saltlength++; + if (saltlength < 0 || saltlength > 19) { + static int warned = 0; + if (warned++ == 1) + fprintf(stderr, + "Note: One or more hashes rejected due to salt length limitation\n"); + return 0; } + + sscanf(&ciphertext[6], "%d", &i); + if (i >= 1 << 16) + return 0; return 1; } @@ -115,7 +107,6 @@ static char *split(char *ciphertext, int { static char out[MAX_CIPHERTEXT_LENGTH + 1]; int i = 0; - for (; ciphertext[i] && i < MAX_CIPHERTEXT_LENGTH; i++) out[i] = ciphertext[i]; out[i] = 0; @@ -127,13 +118,28 @@ static char *split(char *ciphertext, int static char *prepare(char *split_fields[10], struct fmt_main *self) { char *cp; - if (!strncmp(split_fields[1], "$DCC2$", 6) && - valid(split_fields[1], self)) + int i; + if (!strncmp(split_fields[1], "$DCC2$", 6)) { + if (valid(split_fields[1], self)) + return split_fields[1]; + // see if this is a form $DCC2$salt#hash. If so, make it $DCC2$10240#salt#hash and retest (insert 10240# into the line). + cp = mem_alloc(strlen(split_fields[1]) + 7); + sprintf(cp, "$DCC2$10240#%s", &(split_fields[1][6])); + if (valid(cp, self)) { + char *cipher = str_alloc_copy(cp); + MEM_FREE(cp); + return cipher; + } return split_fields[1]; + } if (!split_fields[0]) return split_fields[1]; + // ONLY check, if this string split_fields[1], is ONLY a 32 byte hex string. + for (i = 0; i < 32; i++) + if (atoi16[ARCH_INDEX(split_fields[1][i])] == 0x7F) + return split_fields[1]; cp = mem_alloc(strlen(split_fields[0]) + strlen(split_fields[1]) + 14); - sprintf(cp, "$DCC2$%s#%s", split_fields[0], split_fields[1]); + sprintf(cp, "$DCC2$10240#%s#%s", split_fields[0], split_fields[1]); if (valid(cp, self)) { char *cipher = str_alloc_copy(cp); MEM_FREE(cp); @@ -162,13 +168,25 @@ static void *salt(char *ciphertext) static mscash2_salt salt; char *pos = ciphertext + strlen(mscash2_prefix); int length = 0; + salt.rounds = DEFAULT_ROUNDS; + sscanf(pos, "%d", &salt.rounds); + while (*pos++ != '#'); while (*pos != '#') { if (length == 19) return NULL; salt.salt[length++] = *pos++; } salt.length = length; - //printf("salt len=%d\n",salt.length); +#ifdef _MSCASH2_DEBUG + int i; + printf("salt="); + for (i = 0; i < salt.length; i++) + putchar(salt.salt[i]); + puts(""); + printf("salt len=%d\n", salt.length); + printf("salt rounds=%d\n", salt.rounds); + +#endif return &salt; } @@ -180,11 +198,11 @@ static void set_salt(void *salt) static void set_key(char *key, int index) { #ifdef _MSCASH2_DEBUG - printf("set_key(%d) = [%s]\n",index,key); + printf("set_key(%d) = [%s]\n", index, key); #endif uint8_t length = strlen(key); inbuffer[index].length = length; - memcpy(inbuffer[index].v, key, MIN(length,PLAINTEXT_LENGTH)); + memcpy(inbuffer[index].v, key, MIN(length, PLAINTEXT_LENGTH)); } static char *get_key(int index) @@ -310,51 +328,51 @@ static int cmp_exact(char *source, int c struct fmt_main fmt_cuda_mscash2 = { { - FORMAT_LABEL, - FORMAT_NAME, - ALGORITHM_NAME, - BENCHMARK_COMMENT, - BENCHMARK_LENGTH, - PLAINTEXT_LENGTH, - BINARY_SIZE, - SALT_SIZE, - MIN_KEYS_PER_CRYPT, - MAX_KEYS_PER_CRYPT, - FMT_CASE | FMT_8_BIT| FMT_SPLIT_UNIFIES_CASE | FMT_UNICODE, - tests - }, { - init, - prepare, - valid, - split, - binary, - salt, - { - binary_hash_0, - binary_hash_1, - binary_hash_2, - binary_hash_3, - binary_hash_4, - binary_hash_5, - binary_hash_6 - }, - 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, - get_hash_5, - get_hash_6 - }, - cmp_all, - cmp_one, - cmp_exact + FORMAT_LABEL, + FORMAT_NAME, + ALGORITHM_NAME, + BENCHMARK_COMMENT, + BENCHMARK_LENGTH, + PLAINTEXT_LENGTH, + BINARY_SIZE, + SALT_SIZE, + MIN_KEYS_PER_CRYPT, + MAX_KEYS_PER_CRYPT, + FMT_CASE | FMT_8_BIT | FMT_SPLIT_UNIFIES_CASE | FMT_UNICODE, + tests + },{ + init, + prepare, + valid, + split, + binary, + salt, + { + binary_hash_0, + binary_hash_1, + binary_hash_2, + binary_hash_3, + binary_hash_4, + binary_hash_5, + binary_hash_6 + }, + 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, + get_hash_5, + get_hash_6 + }, + cmp_all, + cmp_one, + cmp_exact } }; diff -urpN magnum-jumbo/src/cuda_mscash_fmt.c magnum-jumbo_mscash-prej7fix//src/cuda_mscash_fmt.c --- magnum-jumbo/src/cuda_mscash_fmt.c 2012-08-08 15:44:55.357683085 +0000 +++ magnum-jumbo_mscash-prej7fix//src/cuda_mscash_fmt.c 2012-08-08 16:10:35.306807966 +0000 @@ -16,7 +16,7 @@ #define FORMAT_LABEL "mscash-cuda" #define FORMAT_NAME "M$ Cache Hash MD4" -#define ALGORITHM_NAME "CUDA, unreliable, may miss guesses" +#define ALGORITHM_NAME "CUDA" #define MAX_CIPHERTEXT_LENGTH (2 + 19*3 + 1 + 32) #define BENCHMARK_COMMENT " len(pass)=8, len(salt)=13" #define BENCHMARK_LENGTH -1 @@ -26,7 +26,18 @@ static mscash_hash *outbuffer; static mscash_salt currentsalt; static struct fmt_tests tests[] = { - {"M$administrator#25fd08fa89795ed54207e6e8442a6ca0", "password"}, + {"M$test1#64cd29e36a8431a2b111378564a10631", "test1"}, + {"M$test1#64cd29e36a8431a2b111378564a10631", "test1"}, + {"M$test1#64cd29e36a8431a2b111378564a10631", "test1"}, + {"176a4c2bd45ac73687676c2f09045353", "", {"root"}}, // nullstring password + {"M$test2#ab60bdb4493822b175486810ac2abe63", "test2"}, + {"M$test3#14dd041848e12fc48c0aa7a416a4a00c", "test3"}, + {"M$test4#b945d24866af4b01a6d89b9d932a153c", "test4"}, + + {"64cd29e36a8431a2b111378564a10631", "test1", {"TEST1"}}, // salt is lowercased before hashing + {"290efa10307e36a79b3eebf2a6b29455", "okolada", {"nineteen_characters"}}, // max salt length + {"ab60bdb4493822b175486810ac2abe63", "test2", {"test2"}}, + {"b945d24866af4b01a6d89b9d932a153c", "test4", {"test4"}}, {NULL} }; @@ -34,19 +45,22 @@ extern void cuda_mscash(mscash_password static void cleanup() { - free(inbuffer); - free(outbuffer); + free(inbuffer); + free(outbuffer); } static void init(struct fmt_main *self) { - //Alocate memory for hashes and passwords - inbuffer=(mscash_password*)calloc(MAX_KEYS_PER_CRYPT, sizeof(mscash_password)); - outbuffer=(mscash_hash*)malloc(sizeof(mscash_hash)*MAX_KEYS_PER_CRYPT); - check_mem_allocation(inbuffer,outbuffer); - atexit(cleanup); - //Initialize CUDA - cuda_init(gpu_id); + //Alocate memory for hashes and passwords + inbuffer = + (mscash_password *) calloc(MAX_KEYS_PER_CRYPT, + sizeof(mscash_password)); + outbuffer = + (mscash_hash *) malloc(sizeof(mscash_hash) * MAX_KEYS_PER_CRYPT); + check_mem_allocation(inbuffer, outbuffer); + atexit(cleanup); + //Initialize CUDA + cuda_init(gpu_id); } static int valid(char *ciphertext, struct fmt_main *self) @@ -76,8 +90,7 @@ static char *split(char *ciphertext, int static char *prepare(char *split_fields[10], struct fmt_main *self) { char *cp; - if (!strncmp(split_fields[1], "M$", 2) && - valid(split_fields[1], self)) + if (!strncmp(split_fields[1], "M$", 2) && valid(split_fields[1], self)) return split_fields[1]; if (!split_fields[0]) return split_fields[1]; @@ -102,7 +115,6 @@ static void *binary(char *ciphertext) binary[i] = SWAP(binary[i]); } return binary; - } static void *salt(char *ciphertext) @@ -110,10 +122,13 @@ static void *salt(char *ciphertext) static mscash_salt salt; char *pos = ciphertext + strlen(mscash_prefix); int length = 0; - while (*pos != '#'){ - if(length == SALT_LENGTH) return NULL; - salt.salt[length++] = *pos++; + while (*pos != '#') { + if (length == SALT_LENGTH) + return NULL; + salt.salt[length++] = *pos++; } + salt.salt[length] = 0; + enc_strlwr(salt.salt); salt.length = length; return &salt; } @@ -127,7 +142,7 @@ static void set_key(char *key, int index { uint8_t length = strlen(key); inbuffer[index].length = length; - memcpy(inbuffer[index].v, key, MIN(length,PLAINTEXT_LENGTH)); + memcpy(inbuffer[index].v, key, MIN(length, PLAINTEXT_LENGTH)); } static char *get_key(int index) @@ -173,6 +188,7 @@ static int binary_hash_5(void *binary) { return ((uint32_t *) binary)[0] & 0xffffff; } + static int binary_hash_6(void *binary) { return ((uint32_t *) binary)[0] & 0x7ffffff; @@ -239,51 +255,51 @@ static int cmp_exact(char *source, int c struct fmt_main fmt_cuda_mscash = { { - FORMAT_LABEL, - FORMAT_NAME, - ALGORITHM_NAME, - BENCHMARK_COMMENT, - BENCHMARK_LENGTH, - PLAINTEXT_LENGTH, - BINARY_SIZE, - SALT_SIZE, - MIN_KEYS_PER_CRYPT, - MAX_KEYS_PER_CRYPT, - FMT_CASE | FMT_8_BIT | FMT_SPLIT_UNIFIES_CASE | FMT_UNICODE, - tests - }, { - init, - prepare, - valid, - split, - binary, - salt, - { - binary_hash_0, - binary_hash_1, - binary_hash_2, - binary_hash_3, - binary_hash_4, - binary_hash_5, - binary_hash_6 - }, - 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, - get_hash_5, - get_hash_6 - }, - cmp_all, - cmp_one, - cmp_exact + FORMAT_LABEL, + FORMAT_NAME, + ALGORITHM_NAME, + BENCHMARK_COMMENT, + BENCHMARK_LENGTH, + PLAINTEXT_LENGTH, + BINARY_SIZE, + SALT_SIZE, + MIN_KEYS_PER_CRYPT, + MAX_KEYS_PER_CRYPT, + FMT_CASE | FMT_8_BIT | FMT_SPLIT_UNIFIES_CASE | FMT_UNICODE, + tests + },{ + init, + prepare, + valid, + split, + binary, + salt, + { + binary_hash_0, + binary_hash_1, + binary_hash_2, + binary_hash_3, + binary_hash_4, + binary_hash_5, + binary_hash_6 + }, + 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, + get_hash_5, + get_hash_6 + }, + cmp_all, + cmp_one, + cmp_exact } };