diff --git a/run/dmg2john.py b/run/dmg2john.py index f3a61a0adb..93264661a9 100755 --- a/run/dmg2john.py +++ b/run/dmg2john.py @@ -92,8 +92,8 @@ def process_file(filename): return data = struct.unpack(v2_header_fmt, data) - (sig, version, enc_iv_size, _, _, _, _, - unk5, uuid, blocksize, datasize, dataoffset, filler1, + (sig, version, enc_iv_size, _, _, data_enc_key_bits, _, + hmac_key_bits, uuid, blocksize, datasize, dataoffset, filler1, kdf_algorithm, kdf_prng_algorithm, kdf_iteration_count, kdf_salt_len, kdf_salt, blob_enc_iv_size, blob_enc_iv, blob_enc_key_bits, blob_enc_algorithm, blob_enc_padding, @@ -111,6 +111,30 @@ def process_file(filename): "is too long!\n" % filename) return + legacy_blob = (blob_enc_algorithm == 17 and blob_enc_key_bits == 192 and + blob_enc_mode == 6 and blob_enc_padding == 7) + aes_blob = (blob_enc_algorithm == 0x80000001 and + blob_enc_key_bits == 192 and blob_enc_mode == 6 and + blob_enc_padding == 7 and encrypted_keyblob_size == 64 and + kdf_salt_len <= 20 and + data_enc_key_bits == 256 and hmac_key_bits == 160 and + kdf_algorithm == 103 and kdf_prng_algorithm == 0) + if not legacy_blob and not aes_blob: + sys.stderr.write("%s uses unsupported blob encryption parameters " \ + "algorithm=%d key_bits=%d mode=%d padding=%d\n" % + (filename, blob_enc_algorithm, blob_enc_key_bits, + blob_enc_mode, blob_enc_padding)) + return + + if aes_blob: + sys.stdout.write("%s:$dmg$3*%d*%s*%d*%d*%s*%d::::%s\n" % + (os.path.basename(filename), kdf_salt_len, + hexlify(kdf_salt)[0:kdf_salt_len*2].decode("ascii"), + blob_enc_key_bits, encrypted_keyblob_size, + hexlify(encrypted_keyblob)[0:encrypted_keyblob_size*2].decode("ascii"), + kdf_iteration_count, filename)) + return + # read starting chunk(s) fd.seek(dataoffset + int(cno * 4096), 0) chunk1 = fd.read(data_size) diff --git a/run/opencl/dmg_kernel.cl b/run/opencl/dmg_kernel.cl index 25c82e65fc..b01bdcee57 100644 --- a/run/opencl/dmg_kernel.cl +++ b/run/opencl/dmg_kernel.cl @@ -24,7 +24,8 @@ typedef struct { uint ivlen; uchar iv[32]; uint32_t encrypted_keyblob_size; - uint8_t encrypted_keyblob[32]; + uint8_t encrypted_keyblob[64]; + uint blob_enc_keybits; uint len_wrapped_aes_key; uchar wrapped_aes_key[296]; uint len_hmac_sha1_key; @@ -136,6 +137,24 @@ INLINE int check_v2hash(const uchar *derived_key, return 0; } +INLINE int check_v3hash(const uchar *derived_key, + MAYBE_CONSTANT dmg_salt *salt, __local aes_local_t *lt) +{ + AES_KEY aes_decrypt_key; aes_decrypt_key.lt = lt; + uchar last_block[16]; + int i; + + AES_set_decrypt_key(derived_key, salt->blob_enc_keybits, &aes_decrypt_key); + memcpy_mcp(last_block, &salt->encrypted_keyblob[48], sizeof(last_block)); + AES_decrypt(last_block, last_block, &aes_decrypt_key); + for (i = 0; i < 16; i++) + last_block[i] ^= salt->encrypted_keyblob[32 + i]; + + return last_block[4] == 'C' && last_block[5] == 'K' && + last_block[6] == 'I' && last_block[7] == 'E' && + !last_block[8] && check_pkcs_pad(last_block, 16, 16) == 9; +} + __kernel void dmg_final_v1(MAYBE_CONSTANT dmg_salt *salt, __global dmg_out *out) @@ -160,3 +179,16 @@ void dmg_final_v2(MAYBE_CONSTANT dmg_salt *salt, out[gid].cracked = check_v2hash((uchar*)dk, salt, <); } + +__kernel +void dmg_final_v3(MAYBE_CONSTANT dmg_salt *salt, + __global dmg_out *out) +{ + __local aes_local_t lt; + uint gid = get_global_id(0); + uint dk[OUTLEN / 4]; + + memcpy_gp(dk, out[gid].dk, OUTLEN); + + out[gid].cracked = check_v3hash((uchar*)dk, salt, <); +} diff --git a/src/dmg2john.c b/src/dmg2john.c index 3ec16b755b..97dd554015 100644 --- a/src/dmg2john.c +++ b/src/dmg2john.c @@ -270,6 +270,7 @@ static void hash_plugin_parse_hash(char *in_filepath) printf("*%u::::%s\n", header.kdf_iteration_count, filename); } else { cencrypted_v2_key_header_pointer header_pointer; + int aes_blob = 0; int password_header_found = 0; if (lseek(fd, 0, SEEK_SET) < 0) { @@ -342,6 +343,30 @@ static void hash_plugin_parse_hash(char *in_filepath) v2_password_header_byteorder_fix(&v2_password_header); + aes_blob = v2_password_header.blob_enc_algo == 0x80000001 && + v2_password_header.blob_enc_keybits == 192 && + v2_password_header.blob_enc_mode == 6 && + v2_password_header.blob_enc_padding == 7 && + v2_password_header.keyblobsize == 64 && + v2_password_header.salt_size <= 20 && + v2_password_header.algorithm == 103 && + v2_password_header.prngalgo == 0 && + header2.keyBits == 256 && + header2.prngkeysize == 160; + + if (!aes_blob && (v2_password_header.blob_enc_algo != 17 || + v2_password_header.blob_enc_keybits != 192 || + v2_password_header.blob_enc_mode != 6 || + v2_password_header.blob_enc_padding != 7)) { + fprintf(stderr, "%s uses unsupported blob encryption parameters " + "algorithm=%u key_bits=%u mode=%u padding=%u\n", filename, + v2_password_header.blob_enc_algo, + v2_password_header.blob_enc_keybits, + v2_password_header.blob_enc_mode, + v2_password_header.blob_enc_padding); + goto bailout; + } + // Allocate the keyblob memory if (v2_password_header.keyblobsize > 1024) { fprintf(stderr, "Unusual keyblobsize found in %s\n", filename); @@ -386,6 +411,18 @@ static void hash_plugin_parse_hash(char *in_filepath) goto bailout; } + if (aes_blob) { + replace(name, ':', ' '); + printf("%s:$dmg$3*%u*", name, v2_password_header.salt_size); + print_hex(v2_password_header.salt, v2_password_header.salt_size); + printf("*%u*%u*", v2_password_header.blob_enc_keybits, + v2_password_header.keyblobsize); + print_hex(v2_password_header.keyblob, v2_password_header.keyblobsize); + printf("*%u::::%s\n", v2_password_header.itercount, filename); + free(v2_password_header.keyblob); + goto bailout; + } + fprintf(stderr, "%s (DMG v%d) successfully parsed, iterations " "count %u\n", name, headerver, v2_password_header.itercount); diff --git a/src/dmg_common_plug.c b/src/dmg_common_plug.c index cfd77455d0..d16c3ba87d 100644 --- a/src/dmg_common_plug.c +++ b/src/dmg_common_plug.c @@ -8,6 +8,8 @@ #include "formats.h" struct fmt_tests dmg_tests[] = { + // Version 2 header using an AES-192-CBC encrypted key blob + {"$dmg$3*20*dda90b3cadb0cfa4e6b17b935c7f415ecb3d4947*192*64*89669cb0994d0bba4d7bdd5fb8551ae0d641a9427cb6903b9546f38c0053c2456799d990a76007f6a7aa3c59fae69889d79df480230d21faaf7a65ab26e45151*400000", "d3fc0n"}, // testimage.AES-256.64k.header_v2.dmg {"$dmg$2*20*fd70ac1e078f01fce55a2e56145a2494446db32a*32*9110b1778f09b1a7000000000000000000000000000000000000000000000000*64*68a32866b0e67515f35dc67c4d6747a8561a9f4f6a6718a894b0a77a47c452471e04ecef9bf56f0d83d1201a509a374e00000000000000000000000000000000*14*8192*70ebe6f1d387e33e3d1093cca2e94c9a32e2c9ba47d461d737d49a7dc1b1f69407b7dbc16f7671689ea4a4641652b3f976b6f1c73c551a0a407d5a335caa169db4a6a25bbd27fbbc38fc71b29ee9b1eae349b0d8a21d57959ecca6bf74bc26ccaee69cfee4999b55374605491af6d0b9066c26995209cd1b71925bcb45a8ef5727a6c20338f08de4357d4cb42cb65ecdc2344a5d7387633c913258ba40699ea5f88804b5e562bf973096337b17b4fc1236d3c8a80b9b48aed63c5a0eae3ae924a883e948f374771bba46923658f225fd2795ce0e795269f589e0ffc81615585e1224cddde654d689a3260e69683c6198bdfcd87507c23cefe36d72f8878cb27bbe5dce868752a7cce067f5a3110f20ebd31ecd53840103e0b2d44385656398edc487bf6d1a5ec3a56af54f9d4254fd20988df41eb85e366f13da1270a3f42c6672ad5faf00fa21e9ba3691bde78ab2c267a142f275467d5b853a107dbf1d75839f0e87b3b4f1d2cec88cc02a26bc4a63aa6836b0c43c5dbb44a832050385a48d46968361ebb053c2416c02458b76c95e50970922556d40b100967340a32824e6b6e44c0c1e0da7ce989d9d5ad91560156" "ed39666cbfbea71f28797a5a7a40e77665612e977ecb8b7fe71d500eafc29d9a0ec1d0ff1723fea7c405bc181ea93c0df42f5bf886eace3cfeee8b0dba52ba8cd2ae009e75d8845264d12dd632ca3236bc1b643437881b270183d2e2bd20808ae73d32bfe88347e33bef4921fcfac9646b74f116be1f04fc353d2222499d5247fa842d0d0f00fc9642ea7524adb65c18fff87b6efd060ec850d7de6f59869387b3d4cc8e38014d52d94ead07d16b8d94327fe5533941497c9be2dd6c04142ba57e29daaeef96d0f2d109522651d797715f4bc5f4cc3fb69fa92623b5ea3e08ff78dc59913993c877f4e2c8964dffd2c8cde6c6b6738da2883505486df5b633aaa8c66acbc2886107f3dd61b1df29f54a13ef27a7d2785c02153375240885e5c54297d88827403320799e05213761549eedc1c159c922087983410d2abadf9ef8ae460d018c278a9ea724f52b866e3d7ff2374496103b5137297100c970d195fca8c1286a8f9d3859ee12c84bdaa4b56ca91e307580b61dbe435ce4021007e4a2a8085976549cf1d195f439bb6e642567f91a0224e98796614d9ea6bfab8f6d13f91b7a80a54e538a1a785cd07b5d7ed2b7e45a0658b5722b5f8844f5139cff3b33ce244946757c020c54c8b5e43324023ed11001201213ffe4829e37135686a8bec1837b35fb234049570868dc5ba9c84cef6890d9ec400a794b1723eb209a60758ba9ae9abd23a7ea9f94fc6b73d29a560e24973c9160f195fbe82376c81dfeec1a7f912a8c22c067a26786a22f0b7db298" diff --git a/src/dmg_fmt_plug.c b/src/dmg_fmt_plug.c index b3377b3c67..a8ab246252 100644 --- a/src/dmg_fmt_plug.c +++ b/src/dmg_fmt_plug.c @@ -134,6 +134,7 @@ static struct custom_salt { unsigned char chunk[8192]; uint32_t encrypted_keyblob_size; uint8_t encrypted_keyblob[128]; + unsigned int blob_enc_keybits; unsigned int len_wrapped_aes_key; unsigned char wrapped_aes_key[296]; unsigned int len_hmac_sha1_key; @@ -177,7 +178,38 @@ static int valid(char *ciphertext, struct fmt_main *self) if ((p = strtokm(ctcopy, "*")) == NULL) goto err; headerver = atoi(p); - if (headerver == 2) { + if (headerver == 3) { + if ((p = strtokm(NULL, "*")) == NULL) /* salt len */ + goto err; + if (!isdec(p)) + goto err; + res = atoi(p); + if (res > 20) + goto err; + if ((p = strtokm(NULL, "*")) == NULL) /* salt */ + goto err; + if (hexlenl(p, &extra) / 2 != res || extra) + goto err; + if ((p = strtokm(NULL, "*")) == NULL) /* blob key bits */ + goto err; + if (!isdec(p) || atoi(p) != 192) + goto err; + if ((p = strtokm(NULL, "*")) == NULL) /* encrypted keyblob size */ + goto err; + if (!isdec(p) || atoi(p) != 64) + goto err; + if ((p = strtokm(NULL, "*")) == NULL) /* encrypted keyblob */ + goto err; + if (hexlenl(p, &extra) != 128 || extra) + goto err; + if ((p = strtokm(NULL, "*")) == NULL) /* iterations */ + goto err; + if (!isdec(p) || !atoi(p)) + goto err; + if (strtokm(NULL, "*") != NULL) + goto err; + } + else if (headerver == 2) { if ((p = strtokm(NULL, "*")) == NULL) /* salt len */ goto err; if (!isdec(p)) @@ -294,7 +326,25 @@ static void *get_salt(char *ciphertext) ctcopy += FORMAT_TAG_LEN; p = strtokm(ctcopy, "*"); cs.headerver = atoi(p); - if (cs.headerver == 2) { + if (cs.headerver == 3) { + p = strtokm(NULL, "*"); + cs.saltlen = atoi(p); + p = strtokm(NULL, "*"); + for (i = 0; i < cs.saltlen; i++) + cs.salt[i] = atoi16[ARCH_INDEX(p[i * 2])] * 16 + + atoi16[ARCH_INDEX(p[i * 2 + 1])]; + p = strtokm(NULL, "*"); + cs.blob_enc_keybits = atoi(p); + p = strtokm(NULL, "*"); + cs.encrypted_keyblob_size = atoi(p); + p = strtokm(NULL, "*"); + for (i = 0; i < cs.encrypted_keyblob_size; i++) + cs.encrypted_keyblob[i] = atoi16[ARCH_INDEX(p[i * 2])] * 16 + + atoi16[ARCH_INDEX(p[i * 2 + 1])]; + p = strtokm(NULL, "*"); + cs.iterations = atoi(p); + } + else if (cs.headerver == 2) { p = strtokm(NULL, "*"); cs.saltlen = atoi(p); p = strtokm(NULL, "*"); @@ -402,7 +452,54 @@ static void hash_plugin_check_hash(int index) unsigned char aes_key_[32]; int j; - if (cur_salt->headerver == 1) { + if (cur_salt->headerver == 3) { +#ifdef SIMD_COEF_32 + unsigned char *derived_key, Derived_key[SSE_GROUP_SZ_SHA1][32]; + int lens[SSE_GROUP_SZ_SHA1], i; + unsigned char *pin[SSE_GROUP_SZ_SHA1]; + union { + uint32_t *pout[SSE_GROUP_SZ_SHA1]; + unsigned char *poutc; + } x; + for (i = 0; i < SSE_GROUP_SZ_SHA1; ++i) { + lens[i] = strlen(saved_key[index+i]); + pin[i] = (unsigned char*)saved_key[index+i]; + x.pout[i] = (uint32_t*)(Derived_key[i]); + } + pbkdf2_sha1_sse((const unsigned char **)pin, lens, cur_salt->salt, + cur_salt->saltlen, + cur_salt->iterations, &(x.poutc), 32, 0); +#else + unsigned char derived_key[32]; + const char *password = saved_key[index]; + pbkdf2_sha1((const unsigned char*)password, strlen(password), + cur_salt->salt, cur_salt->saltlen, cur_salt->iterations, + derived_key, 32, 0); +#endif + j = 0; +#ifdef SIMD_COEF_32 + for (j = 0; j < SSE_GROUP_SZ_SHA1; ++j) { + derived_key = Derived_key[j]; +#endif + AES_KEY aes_decrypt_key; + unsigned char last_block[16]; + int i; + + AES_set_decrypt_key(derived_key, cur_salt->blob_enc_keybits, + &aes_decrypt_key); + AES_decrypt(&cur_salt->encrypted_keyblob[48], last_block, + &aes_decrypt_key); + for (i = 0; i < 16; i++) + last_block[i] ^= cur_salt->encrypted_keyblob[32 + i]; + + if (!memcmp(&last_block[4], "CKIE", 4) && !last_block[8] && + check_pkcs_pad(last_block, 16, 16) == 9) + cracked[index+j] = 1; +#ifdef SIMD_COEF_32 + } +#endif + } + else if (cur_salt->headerver == 1) { #ifdef SIMD_COEF_32 unsigned char *derived_key, Derived_key[SSE_GROUP_SZ_SHA1][32]; int lens[SSE_GROUP_SZ_SHA1], i; diff --git a/src/opencl_dmg_fmt_plug.c b/src/opencl_dmg_fmt_plug.c index 09d549b653..e4a7d41ea9 100644 --- a/src/opencl_dmg_fmt_plug.c +++ b/src/opencl_dmg_fmt_plug.c @@ -52,7 +52,8 @@ typedef struct { unsigned int ivlen; unsigned char iv[32]; uint32_t encrypted_keyblob_size; - uint8_t encrypted_keyblob[32]; + uint8_t encrypted_keyblob[64]; + unsigned int blob_enc_keybits; unsigned int len_wrapped_aes_key; unsigned char wrapped_aes_key[296]; unsigned int len_hmac_sha1_key; @@ -78,7 +79,7 @@ static cl_mem mem_in, mem_out, mem_salt, mem_state; static int new_keys; static struct fmt_main *self; -static cl_kernel pbkdf2_init, pbkdf2_loop, pbkdf2_final, dmg_final[3]; +static cl_kernel pbkdf2_init, pbkdf2_loop, pbkdf2_final, dmg_final[4]; /* * HASH_LOOPS is ideally made by factors of (iteration count - 1) and should @@ -108,6 +109,7 @@ static size_t get_task_max_work_group_size() s = MIN(s, autotune_get_task_max_work_group_size(FALSE, 0, pbkdf2_final)); s = MIN(s, autotune_get_task_max_work_group_size(FALSE, 0, dmg_final[1])); s = MIN(s, autotune_get_task_max_work_group_size(FALSE, 0, dmg_final[2])); + s = MIN(s, autotune_get_task_max_work_group_size(FALSE, 0, dmg_final[3])); return s; } @@ -151,6 +153,9 @@ static void create_clobj(size_t gws, struct fmt_main *self) HANDLE_CLERROR(clSetKernelArg(dmg_final[2], 0, sizeof(mem_salt), &mem_salt), "Error while setting mem_salt kernel argument"); HANDLE_CLERROR(clSetKernelArg(dmg_final[2], 1, sizeof(mem_out), &mem_out), "Error while setting mem_out kernel argument"); + + HANDLE_CLERROR(clSetKernelArg(dmg_final[3], 0, sizeof(mem_salt), &mem_salt), "Error while setting mem_salt kernel argument"); + HANDLE_CLERROR(clSetKernelArg(dmg_final[3], 1, sizeof(mem_out), &mem_out), "Error while setting mem_out kernel argument"); } static void release_clobj(void) @@ -177,6 +182,7 @@ static void done(void) HANDLE_CLERROR(clReleaseKernel(pbkdf2_final), "Release kernel"); HANDLE_CLERROR(clReleaseKernel(dmg_final[1]), "Release kernel"); HANDLE_CLERROR(clReleaseKernel(dmg_final[2]), "Release kernel"); + HANDLE_CLERROR(clReleaseKernel(dmg_final[3]), "Release kernel"); HANDLE_CLERROR(clReleaseProgram(program[gpu_id]), "Release Program"); program[gpu_id] = NULL; @@ -223,6 +229,8 @@ static void reset(struct db_main *db) HANDLE_CLERROR(ret_code, "Error creating kernel"); dmg_final[2] = clCreateKernel(program[gpu_id], "dmg_final_v2", &ret_code); HANDLE_CLERROR(ret_code, "Error creating kernel"); + dmg_final[3] = clCreateKernel(program[gpu_id], "dmg_final_v3", &ret_code); + HANDLE_CLERROR(ret_code, "Error creating kernel"); } // FIXME: Share in opencl_autotune.h @@ -253,7 +261,38 @@ static int valid(char *ciphertext, struct fmt_main *self) if ((p = strtokm(ctcopy, "*")) == NULL) goto err; headerver = atoi(p); - if (headerver == 2) { + if (headerver == 3) { + if ((p = strtokm(NULL, "*")) == NULL) /* salt len */ + goto err; + if (!isdec(p)) + goto err; + res = atoi(p); + if (res > 20) + goto err; + if ((p = strtokm(NULL, "*")) == NULL) /* salt */ + goto err; + if (hexlenl(p, &extra) != res*2 || extra) + goto err; + if ((p = strtokm(NULL, "*")) == NULL) /* blob key bits */ + goto err; + if (!isdec(p) || atoi(p) != 192) + goto err; + if ((p = strtokm(NULL, "*")) == NULL) /* encrypted keyblob size */ + goto err; + if (!isdec(p) || atoi(p) != 64) + goto err; + if ((p = strtokm(NULL, "*")) == NULL) /* encrypted keyblob */ + goto err; + if (hexlenl(p, &extra) != 128 || extra) + goto err; + if ((p = strtokm(NULL, "*")) == NULL) /* iterations */ + goto err; + if (!isdec(p) || !atoi(p)) + goto err; + if (strtokm(NULL, "*") != NULL) + goto err; + } + else if (headerver == 2) { if ((p = strtokm(NULL, "*")) == NULL) /* salt len */ goto err; if (!isdec(p)) @@ -370,7 +409,25 @@ static void *get_salt(char *ciphertext) ctcopy += FORMAT_TAG_LEN; p = strtokm(ctcopy, "*"); cs.headerver = atoi(p); - if (cs.headerver == 2) { + if (cs.headerver == 3) { + p = strtokm(NULL, "*"); + cs.pbkdf2.length = atoi(p); + p = strtokm(NULL, "*"); + for (i = 0; i < cs.pbkdf2.length; i++) + cs.pbkdf2.salt[i] = atoi16[ARCH_INDEX(p[i * 2])] * 16 + + atoi16[ARCH_INDEX(p[i * 2 + 1])]; + p = strtokm(NULL, "*"); + cs.blob_enc_keybits = atoi(p); + p = strtokm(NULL, "*"); + cs.encrypted_keyblob_size = atoi(p); + p = strtokm(NULL, "*"); + for (i = 0; i < cs.encrypted_keyblob_size; i++) + cs.encrypted_keyblob[i] = atoi16[ARCH_INDEX(p[i * 2])] * 16 + + atoi16[ARCH_INDEX(p[i * 2 + 1])]; + p = strtokm(NULL, "*"); + cs.pbkdf2.iterations = atoi(p); + } + else if (cs.headerver == 2) { p = strtokm(NULL, "*"); cs.pbkdf2.length = atoi(p); p = strtokm(NULL, "*");