diff --git a/OpenCL/inc_truecrypt_keyboard.cl b/OpenCL/inc_truecrypt_keyboard.cl index 8d6fb7880..9afd8c717 100644 --- a/OpenCL/inc_truecrypt_keyboard.cl +++ b/OpenCL/inc_truecrypt_keyboard.cl @@ -1,22 +1,143 @@ -DECLSPEC void keyboard_map (u32 w[4], __local u32 *s_keyboard_layout) +DECLSPEC int find_map (const u32 search, const int search_len, __local kb_layout_map_t *s_kb_layout_map, const int kb_layout_map_cnt) { - w[0] = (s_keyboard_layout[(w[0] >> 0) & 0xff] << 0) - | (s_keyboard_layout[(w[0] >> 8) & 0xff] << 8) - | (s_keyboard_layout[(w[0] >> 16) & 0xff] << 16) - | (s_keyboard_layout[(w[0] >> 24) & 0xff] << 24); - - w[1] = (s_keyboard_layout[(w[1] >> 0) & 0xff] << 0) - | (s_keyboard_layout[(w[1] >> 8) & 0xff] << 8) - | (s_keyboard_layout[(w[1] >> 16) & 0xff] << 16) - | (s_keyboard_layout[(w[1] >> 24) & 0xff] << 24); - - w[2] = (s_keyboard_layout[(w[2] >> 0) & 0xff] << 0) - | (s_keyboard_layout[(w[2] >> 8) & 0xff] << 8) - | (s_keyboard_layout[(w[2] >> 16) & 0xff] << 16) - | (s_keyboard_layout[(w[2] >> 24) & 0xff] << 24); - - w[3] = (s_keyboard_layout[(w[3] >> 0) & 0xff] << 0) - | (s_keyboard_layout[(w[3] >> 8) & 0xff] << 8) - | (s_keyboard_layout[(w[3] >> 16) & 0xff] << 16) - | (s_keyboard_layout[(w[3] >> 24) & 0xff] << 24); + for (int idx = 0; idx < kb_layout_map_cnt; idx++) + { + const u32 src_char = s_kb_layout_map[idx].src_char; + const int src_len = s_kb_layout_map[idx].src_len; + + if (src_len == search_len) + { + const u32 mask = 0xffffffff >> ((4 - search_len) * 8); + + if ((src_char & mask) == (search & mask)) return idx; + } + } + + return -1; +} + +DECLSPEC int keyboard_map (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int pw_len, __local kb_layout_map_t *s_kb_layout_map, const int kb_layout_map_cnt) +{ + u32 out_buf[16] = { 0 }; + + u8 *out_ptr = (u8 *) out_buf; + + int out_len = 0; + + // TC/VC passwords are limited to 64 + + u32 w[16]; + + w[ 0] = w0[0]; + w[ 1] = w0[1]; + w[ 2] = w0[2]; + w[ 3] = w0[3]; + w[ 4] = w1[0]; + w[ 5] = w1[1]; + w[ 6] = w1[2]; + w[ 7] = w1[3]; + w[ 8] = w2[0]; + w[ 9] = w2[1]; + w[10] = w2[2]; + w[11] = w2[3]; + w[12] = w3[0]; + w[13] = w3[1]; + w[14] = w3[2]; + w[15] = w3[3]; + + u8 *w_ptr = (u8 *) w; + + int pw_pos = 0; + + while (pw_pos < pw_len) + { + u32 src0 = 0; + u32 src1 = 0; + u32 src2 = 0; + u32 src3 = 0; + + #define MIN(a,b) (((a) < (b)) ? (a) : (b)) + + const int rem = MIN (pw_len - pw_pos, 4); + + #undef MIN + + if (rem > 0) src0 = w_ptr[pw_pos + 0]; + if (rem > 1) src1 = w_ptr[pw_pos + 1]; + if (rem > 2) src2 = w_ptr[pw_pos + 2]; + if (rem > 3) src3 = w_ptr[pw_pos + 3]; + + const u32 src = (src0 << 0) + | (src1 << 8) + | (src2 << 16) + | (src3 << 24); + + int src_len; + + for (src_len = rem; src_len > 0; src_len--) + { + const int idx = find_map (src, src_len, s_kb_layout_map, kb_layout_map_cnt); + + if (idx == -1) continue; + + u32 dst_char = s_kb_layout_map[idx].dst_char; + int dst_len = s_kb_layout_map[idx].dst_len; + + switch (dst_len) + { + case 1: + out_ptr[out_len++] = (dst_char >> 0) & 0xff; + break; + case 2: + out_ptr[out_len++] = (dst_char >> 0) & 0xff; + out_ptr[out_len++] = (dst_char >> 8) & 0xff; + break; + case 3: + out_ptr[out_len++] = (dst_char >> 0) & 0xff; + out_ptr[out_len++] = (dst_char >> 8) & 0xff; + out_ptr[out_len++] = (dst_char >> 16) & 0xff; + break; + case 4: + out_ptr[out_len++] = (dst_char >> 0) & 0xff; + out_ptr[out_len++] = (dst_char >> 8) & 0xff; + out_ptr[out_len++] = (dst_char >> 16) & 0xff; + out_ptr[out_len++] = (dst_char >> 24) & 0xff; + break; + } + + pw_pos += src_len; + + break; + } + + // not matched, keep original + + if (src_len == 0) + { + out_ptr[out_len] = w_ptr[pw_pos]; + + out_len++; + + pw_pos++; + } + } + + w0[0] = out_buf[ 0]; + w0[1] = out_buf[ 1]; + w0[2] = out_buf[ 2]; + w0[3] = out_buf[ 3]; + w1[0] = out_buf[ 4]; + w1[1] = out_buf[ 5]; + w1[2] = out_buf[ 6]; + w1[3] = out_buf[ 7]; + w2[0] = out_buf[ 8]; + w2[1] = out_buf[ 9]; + w2[2] = out_buf[10]; + w2[3] = out_buf[11]; + w3[0] = out_buf[12]; + w3[1] = out_buf[13]; + w3[2] = out_buf[14]; + w3[3] = out_buf[15]; + + return out_len; } diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl index 1a692b7db..4dcf86ec1 100644 --- a/OpenCL/inc_types.cl +++ b/OpenCL/inc_types.cl @@ -1303,14 +1303,25 @@ typedef struct krb5asrep } krb5asrep_t; +typedef struct kb_layout_map +{ + u32 src_char; + int src_len; + u32 dst_char; + int dst_len; + +} kb_layout_map_t; + typedef struct tc { u32 salt_buf[32]; u32 data_buf[112]; u32 keyfile_buf[16]; - u32 keyboard_layout[256]; u32 signature; + kb_layout_map_t kb_layout_map[256]; + int kb_layout_map_cnt; + } tc_t; typedef struct pbkdf2_md5 diff --git a/OpenCL/m06211-pure.cl b/OpenCL/m06211-pure.cl index 2c53edcd6..04f0fbf16 100644 --- a/OpenCL/m06211-pure.cl +++ b/OpenCL/m06211-pure.cl @@ -68,11 +68,13 @@ __kernel void m06211_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) * keyboard layout shared */ - __local u32 s_keyboard_layout[256]; + const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt; + + __local kb_layout_map_t s_kb_layout_map[256]; for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) { - s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i]; } barrier (CLK_LOCAL_MEM_FENCE); @@ -105,10 +107,9 @@ __kernel void m06211_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) w3[2] = pws[gid].i[14]; w3[3] = pws[gid].i[15]; - keyboard_map (w0, s_keyboard_layout); - keyboard_map (w1, s_keyboard_layout); - keyboard_map (w2, s_keyboard_layout); - keyboard_map (w3, s_keyboard_layout); + const u32 pw_len = pws[gid].pw_len; + + keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_cnt); w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); diff --git a/OpenCL/m06212-pure.cl b/OpenCL/m06212-pure.cl index 25ad6bca8..b316eba04 100644 --- a/OpenCL/m06212-pure.cl +++ b/OpenCL/m06212-pure.cl @@ -68,11 +68,13 @@ __kernel void m06212_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) * keyboard layout shared */ - __local u32 s_keyboard_layout[256]; + const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt; + + __local kb_layout_map_t s_kb_layout_map[256]; for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) { - s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i]; } barrier (CLK_LOCAL_MEM_FENCE); @@ -105,10 +107,9 @@ __kernel void m06212_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) w3[2] = pws[gid].i[14]; w3[3] = pws[gid].i[15]; - keyboard_map (w0, s_keyboard_layout); - keyboard_map (w1, s_keyboard_layout); - keyboard_map (w2, s_keyboard_layout); - keyboard_map (w3, s_keyboard_layout); + const u32 pw_len = pws[gid].pw_len; + + keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_cnt); w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); diff --git a/OpenCL/m06213-pure.cl b/OpenCL/m06213-pure.cl index a74a64e7b..e7eeb146a 100644 --- a/OpenCL/m06213-pure.cl +++ b/OpenCL/m06213-pure.cl @@ -68,11 +68,13 @@ __kernel void m06213_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) * keyboard layout shared */ - __local u32 s_keyboard_layout[256]; + const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt; + + __local kb_layout_map_t s_kb_layout_map[256]; for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) { - s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i]; } barrier (CLK_LOCAL_MEM_FENCE); @@ -105,10 +107,9 @@ __kernel void m06213_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) w3[2] = pws[gid].i[14]; w3[3] = pws[gid].i[15]; - keyboard_map (w0, s_keyboard_layout); - keyboard_map (w1, s_keyboard_layout); - keyboard_map (w2, s_keyboard_layout); - keyboard_map (w3, s_keyboard_layout); + const u32 pw_len = pws[gid].pw_len; + + keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_cnt); w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); diff --git a/OpenCL/m06221-pure.cl b/OpenCL/m06221-pure.cl index e9b03f5bf..0ddf34407 100644 --- a/OpenCL/m06221-pure.cl +++ b/OpenCL/m06221-pure.cl @@ -92,11 +92,13 @@ __kernel void m06221_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t)) * keyboard layout shared */ - __local u32 s_keyboard_layout[256]; + const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt; + + __local kb_layout_map_t s_kb_layout_map[256]; for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) { - s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i]; } barrier (CLK_LOCAL_MEM_FENCE); @@ -149,14 +151,9 @@ __kernel void m06221_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t)) w7[2] = pws[gid].i[30]; w7[3] = pws[gid].i[31]; - keyboard_map (w0, s_keyboard_layout); - keyboard_map (w1, s_keyboard_layout); - keyboard_map (w2, s_keyboard_layout); - keyboard_map (w3, s_keyboard_layout); - keyboard_map (w4, s_keyboard_layout); - keyboard_map (w5, s_keyboard_layout); - keyboard_map (w6, s_keyboard_layout); - keyboard_map (w7, s_keyboard_layout); + const u32 pw_len = pws[gid].pw_len; + + keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_cnt); w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); diff --git a/OpenCL/m06222-pure.cl b/OpenCL/m06222-pure.cl index bc865cfdc..d6c2ff8e6 100644 --- a/OpenCL/m06222-pure.cl +++ b/OpenCL/m06222-pure.cl @@ -92,11 +92,13 @@ __kernel void m06222_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t)) * keyboard layout shared */ - __local u32 s_keyboard_layout[256]; + const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt; + + __local kb_layout_map_t s_kb_layout_map[256]; for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) { - s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i]; } barrier (CLK_LOCAL_MEM_FENCE); @@ -149,14 +151,9 @@ __kernel void m06222_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t)) w7[2] = pws[gid].i[30]; w7[3] = pws[gid].i[31]; - keyboard_map (w0, s_keyboard_layout); - keyboard_map (w1, s_keyboard_layout); - keyboard_map (w2, s_keyboard_layout); - keyboard_map (w3, s_keyboard_layout); - keyboard_map (w4, s_keyboard_layout); - keyboard_map (w5, s_keyboard_layout); - keyboard_map (w6, s_keyboard_layout); - keyboard_map (w7, s_keyboard_layout); + const u32 pw_len = pws[gid].pw_len; + + keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_cnt); w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); diff --git a/OpenCL/m06223-pure.cl b/OpenCL/m06223-pure.cl index 2576f7f79..0c0bc2df8 100644 --- a/OpenCL/m06223-pure.cl +++ b/OpenCL/m06223-pure.cl @@ -92,11 +92,13 @@ __kernel void m06223_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t)) * keyboard layout shared */ - __local u32 s_keyboard_layout[256]; + const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt; + + __local kb_layout_map_t s_kb_layout_map[256]; for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) { - s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i]; } barrier (CLK_LOCAL_MEM_FENCE); @@ -149,14 +151,9 @@ __kernel void m06223_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t)) w7[2] = pws[gid].i[30]; w7[3] = pws[gid].i[31]; - keyboard_map (w0, s_keyboard_layout); - keyboard_map (w1, s_keyboard_layout); - keyboard_map (w2, s_keyboard_layout); - keyboard_map (w3, s_keyboard_layout); - keyboard_map (w4, s_keyboard_layout); - keyboard_map (w5, s_keyboard_layout); - keyboard_map (w6, s_keyboard_layout); - keyboard_map (w7, s_keyboard_layout); + const u32 pw_len = pws[gid].pw_len; + + keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_cnt); w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); diff --git a/OpenCL/m06231-pure.cl b/OpenCL/m06231-pure.cl index 5df914052..334dad297 100644 --- a/OpenCL/m06231-pure.cl +++ b/OpenCL/m06231-pure.cl @@ -130,11 +130,13 @@ __kernel void m06231_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) * keyboard layout shared */ - __local u32 s_keyboard_layout[256]; + const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt; + + __local kb_layout_map_t s_kb_layout_map[256]; for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) { - s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i]; } __local u32 s_Ch[8][256]; @@ -191,10 +193,9 @@ __kernel void m06231_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) w3[2] = pws[gid].i[14]; w3[3] = pws[gid].i[15]; - keyboard_map (w0, s_keyboard_layout); - keyboard_map (w1, s_keyboard_layout); - keyboard_map (w2, s_keyboard_layout); - keyboard_map (w3, s_keyboard_layout); + const u32 pw_len = pws[gid].pw_len; + + keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_cnt); w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); diff --git a/OpenCL/m06232-pure.cl b/OpenCL/m06232-pure.cl index f09c1333c..d2a411a8c 100644 --- a/OpenCL/m06232-pure.cl +++ b/OpenCL/m06232-pure.cl @@ -130,11 +130,13 @@ __kernel void m06232_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) * keyboard layout shared */ - __local u32 s_keyboard_layout[256]; + const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt; + + __local kb_layout_map_t s_kb_layout_map[256]; for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) { - s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i]; } __local u32 s_Ch[8][256]; @@ -191,10 +193,9 @@ __kernel void m06232_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) w3[2] = pws[gid].i[14]; w3[3] = pws[gid].i[15]; - keyboard_map (w0, s_keyboard_layout); - keyboard_map (w1, s_keyboard_layout); - keyboard_map (w2, s_keyboard_layout); - keyboard_map (w3, s_keyboard_layout); + const u32 pw_len = pws[gid].pw_len; + + keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_cnt); w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); diff --git a/OpenCL/m06233-pure.cl b/OpenCL/m06233-pure.cl index c6ce43cbe..5d8455a59 100644 --- a/OpenCL/m06233-pure.cl +++ b/OpenCL/m06233-pure.cl @@ -130,11 +130,13 @@ __kernel void m06233_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) * keyboard layout shared */ - __local u32 s_keyboard_layout[256]; + const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt; + + __local kb_layout_map_t s_kb_layout_map[256]; for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) { - s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i]; } __local u32 s_Ch[8][256]; @@ -191,10 +193,9 @@ __kernel void m06233_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) w3[2] = pws[gid].i[14]; w3[3] = pws[gid].i[15]; - keyboard_map (w0, s_keyboard_layout); - keyboard_map (w1, s_keyboard_layout); - keyboard_map (w2, s_keyboard_layout); - keyboard_map (w3, s_keyboard_layout); + const u32 pw_len = pws[gid].pw_len; + + keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_cnt); w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); diff --git a/OpenCL/m13751-pure.cl b/OpenCL/m13751-pure.cl index 072548611..a249e85c0 100644 --- a/OpenCL/m13751-pure.cl +++ b/OpenCL/m13751-pure.cl @@ -76,11 +76,13 @@ __kernel void m13751_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) * keyboard layout shared */ - __local u32 s_keyboard_layout[256]; + const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt; + + __local kb_layout_map_t s_kb_layout_map[256]; for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) { - s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i]; } barrier (CLK_LOCAL_MEM_FENCE); @@ -113,10 +115,9 @@ __kernel void m13751_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) w3[2] = pws[gid].i[14]; w3[3] = pws[gid].i[15]; - keyboard_map (w0, s_keyboard_layout); - keyboard_map (w1, s_keyboard_layout); - keyboard_map (w2, s_keyboard_layout); - keyboard_map (w3, s_keyboard_layout); + const u32 pw_len = pws[gid].pw_len; + + keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_cnt); w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); diff --git a/OpenCL/m13752-pure.cl b/OpenCL/m13752-pure.cl index 3589d36a6..f378111f7 100644 --- a/OpenCL/m13752-pure.cl +++ b/OpenCL/m13752-pure.cl @@ -76,11 +76,13 @@ __kernel void m13752_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) * keyboard layout shared */ - __local u32 s_keyboard_layout[256]; + const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt; + + __local kb_layout_map_t s_kb_layout_map[256]; for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) { - s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i]; } barrier (CLK_LOCAL_MEM_FENCE); @@ -113,10 +115,9 @@ __kernel void m13752_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) w3[2] = pws[gid].i[14]; w3[3] = pws[gid].i[15]; - keyboard_map (w0, s_keyboard_layout); - keyboard_map (w1, s_keyboard_layout); - keyboard_map (w2, s_keyboard_layout); - keyboard_map (w3, s_keyboard_layout); + const u32 pw_len = pws[gid].pw_len; + + keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_cnt); w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); diff --git a/OpenCL/m13753-pure.cl b/OpenCL/m13753-pure.cl index f5b2d7b9e..a8ce2fd5c 100644 --- a/OpenCL/m13753-pure.cl +++ b/OpenCL/m13753-pure.cl @@ -76,11 +76,13 @@ __kernel void m13753_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) * keyboard layout shared */ - __local u32 s_keyboard_layout[256]; + const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt; + + __local kb_layout_map_t s_kb_layout_map[256]; for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) { - s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i]; } barrier (CLK_LOCAL_MEM_FENCE); @@ -113,10 +115,9 @@ __kernel void m13753_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) w3[2] = pws[gid].i[14]; w3[3] = pws[gid].i[15]; - keyboard_map (w0, s_keyboard_layout); - keyboard_map (w1, s_keyboard_layout); - keyboard_map (w2, s_keyboard_layout); - keyboard_map (w3, s_keyboard_layout); + const u32 pw_len = pws[gid].pw_len; + + keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_cnt); w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); diff --git a/OpenCL/m13771-pure.cl b/OpenCL/m13771-pure.cl index fbd454bd5..b48e9b495 100644 --- a/OpenCL/m13771-pure.cl +++ b/OpenCL/m13771-pure.cl @@ -115,11 +115,13 @@ __kernel void m13771_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t)) const u64 lid = get_local_id (0); const u64 lsz = get_local_size (0); - __local u32 s_keyboard_layout[256]; + const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt; + + __local kb_layout_map_t s_kb_layout_map[256]; for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) { - s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i]; } #ifdef REAL_SHM @@ -174,10 +176,9 @@ __kernel void m13771_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t)) w3[2] = pws[gid].i[14]; w3[3] = pws[gid].i[15]; - keyboard_map (w0, s_keyboard_layout); - keyboard_map (w1, s_keyboard_layout); - keyboard_map (w2, s_keyboard_layout); - keyboard_map (w3, s_keyboard_layout); + const u32 pw_len = pws[gid].pw_len; + + keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_cnt); w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); diff --git a/OpenCL/m13772-pure.cl b/OpenCL/m13772-pure.cl index 34aa8676e..79632479b 100644 --- a/OpenCL/m13772-pure.cl +++ b/OpenCL/m13772-pure.cl @@ -115,11 +115,13 @@ __kernel void m13772_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t)) const u64 lid = get_local_id (0); const u64 lsz = get_local_size (0); - __local u32 s_keyboard_layout[256]; + const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt; + + __local kb_layout_map_t s_kb_layout_map[256]; for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) { - s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i]; } #ifdef REAL_SHM @@ -174,10 +176,9 @@ __kernel void m13772_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t)) w3[2] = pws[gid].i[14]; w3[3] = pws[gid].i[15]; - keyboard_map (w0, s_keyboard_layout); - keyboard_map (w1, s_keyboard_layout); - keyboard_map (w2, s_keyboard_layout); - keyboard_map (w3, s_keyboard_layout); + const u32 pw_len = pws[gid].pw_len; + + keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_cnt); w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); diff --git a/OpenCL/m13773-pure.cl b/OpenCL/m13773-pure.cl index 8e09a601b..8f6155359 100644 --- a/OpenCL/m13773-pure.cl +++ b/OpenCL/m13773-pure.cl @@ -115,11 +115,13 @@ __kernel void m13773_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t)) const u64 lid = get_local_id (0); const u64 lsz = get_local_size (0); - __local u32 s_keyboard_layout[256]; + const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt; + + __local kb_layout_map_t s_kb_layout_map[256]; for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) { - s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i]; } #ifdef REAL_SHM @@ -174,10 +176,9 @@ __kernel void m13773_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t)) w3[2] = pws[gid].i[14]; w3[3] = pws[gid].i[15]; - keyboard_map (w0, s_keyboard_layout); - keyboard_map (w1, s_keyboard_layout); - keyboard_map (w2, s_keyboard_layout); - keyboard_map (w3, s_keyboard_layout); + const u32 pw_len = pws[gid].pw_len; + + keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_cnt); w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); diff --git a/include/interface.h b/include/interface.h index e4daabf8e..452a9f19a 100644 --- a/include/interface.h +++ b/include/interface.h @@ -317,14 +317,25 @@ typedef struct keepass } keepass_t; +typedef struct kb_layout_map +{ + u32 src_char; + int src_len; + u32 dst_char; + int dst_len; + +} kb_layout_map_t; + typedef struct tc { u32 salt_buf[32]; u32 data_buf[112]; u32 keyfile_buf[16]; - u32 keyboard_layout[256]; u32 signature; + kb_layout_map_t kb_layout_map[256]; + int kb_layout_map_cnt; + } tc_t; typedef struct pbkdf2_md5 diff --git a/src/hashes.c b/src/hashes.c index a72ae6889..fd2176e2e 100644 --- a/src/hashes.c +++ b/src/hashes.c @@ -130,6 +130,129 @@ int sort_by_hash_no_salt (const void *v1, const void *v2, void *v3) return sort_by_digest_p0p1 (d1, d2, v3); } +int find_map (const u32 search, const int search_len, kb_layout_map_t *s_kb_layout_map, const int kb_layout_map_cnt) +{ + for (int idx = 0; idx < kb_layout_map_cnt; idx++) + { + const u32 src_char = s_kb_layout_map[idx].src_char; + const int src_len = s_kb_layout_map[idx].src_len; + + if (src_len == search_len) + { + const u32 mask = 0xffffffff >> ((4 - search_len) * 8); + + if ((src_char & mask) == (search & mask)) return idx; + } + } + + return -1; +} + +int keyboard_map (u32 plain_buf[64], const int plain_len, kb_layout_map_t *s_kb_layout_map, const int kb_layout_map_cnt) +{ + u32 out_buf[16] = { 0 }; + + u8 *out_ptr = (u8 *) out_buf; + + int out_len = 0; + + u8 *plain_ptr = (u8 *) plain_buf; + + int plain_pos = 0; + + while (plain_pos < plain_len) + { + u32 src0 = 0; + u32 src1 = 0; + u32 src2 = 0; + u32 src3 = 0; + + #define MIN(a,b) (((a) < (b)) ? (a) : (b)) + + const int rem = MIN (plain_len - plain_pos, 4); + + #undef MIN + + if (rem > 0) src0 = plain_ptr[plain_pos + 0]; + if (rem > 1) src1 = plain_ptr[plain_pos + 1]; + if (rem > 2) src2 = plain_ptr[plain_pos + 2]; + if (rem > 3) src3 = plain_ptr[plain_pos + 3]; + + const u32 src = (src0 << 0) + | (src1 << 8) + | (src2 << 16) + | (src3 << 24); + + int src_len; + + for (src_len = rem; src_len > 0; src_len--) + { + const int idx = find_map (src, src_len, s_kb_layout_map, kb_layout_map_cnt); + + if (idx == -1) continue; + + u32 dst_char = s_kb_layout_map[idx].dst_char; + int dst_len = s_kb_layout_map[idx].dst_len; + + switch (dst_len) + { + case 1: + out_ptr[out_len++] = (dst_char >> 0) & 0xff; + break; + case 2: + out_ptr[out_len++] = (dst_char >> 0) & 0xff; + out_ptr[out_len++] = (dst_char >> 8) & 0xff; + break; + case 3: + out_ptr[out_len++] = (dst_char >> 0) & 0xff; + out_ptr[out_len++] = (dst_char >> 8) & 0xff; + out_ptr[out_len++] = (dst_char >> 16) & 0xff; + break; + case 4: + out_ptr[out_len++] = (dst_char >> 0) & 0xff; + out_ptr[out_len++] = (dst_char >> 8) & 0xff; + out_ptr[out_len++] = (dst_char >> 16) & 0xff; + out_ptr[out_len++] = (dst_char >> 24) & 0xff; + break; + } + + plain_pos += src_len; + + break; + } + + // not matched, keep original + + if (src_len == 0) + { + out_ptr[out_len] = plain_ptr[plain_pos]; + + out_len++; + + plain_pos++; + } + } + + plain_buf[ 0] = out_buf[ 0]; + plain_buf[ 1] = out_buf[ 1]; + plain_buf[ 2] = out_buf[ 2]; + plain_buf[ 3] = out_buf[ 3]; + plain_buf[ 4] = out_buf[ 4]; + plain_buf[ 5] = out_buf[ 5]; + plain_buf[ 6] = out_buf[ 6]; + plain_buf[ 7] = out_buf[ 7]; + plain_buf[ 8] = out_buf[ 8]; + plain_buf[ 9] = out_buf[ 9]; + plain_buf[10] = out_buf[10]; + plain_buf[11] = out_buf[11]; + plain_buf[12] = out_buf[12]; + plain_buf[13] = out_buf[13]; + plain_buf[14] = out_buf[14]; + plain_buf[15] = out_buf[15]; + + return out_len; +} + int save_hash (hashcat_ctx_t *hashcat_ctx) { hashes_t *hashes = hashcat_ctx->hashes; @@ -317,10 +440,7 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl { tc_t *tc = (tc_t *) hashes->esalts_buf; - for (int i = 0; i < plain_len; i++) - { - plain_ptr[i] = (u8) tc->keyboard_layout[plain_ptr[i]]; - } + plain_len = keyboard_map (plain_buf, plain_len, tc->kb_layout_map, tc->kb_layout_map_cnt); } // crackpos diff --git a/src/interface.c b/src/interface.c index 2eddcf892..dcc0748e1 100644 --- a/src/interface.c +++ b/src/interface.c @@ -2628,7 +2628,15 @@ static int input_tokenizer (u8 *input_buf, int input_len, token_t *token) return PARSER_OK; } -static bool initialize_keyboard_layout (hashcat_ctx_t *hashcat_ctx, const char *filename, u32 *keyboard_layout) +static int sort_by_src_len (const void *p1, const void *p2) +{ + const kb_layout_map_t *k1 = (const kb_layout_map_t *) p1; + const kb_layout_map_t *k2 = (const kb_layout_map_t *) p2; + + return k1->src_len < k2->src_len; +} + +static bool initialize_keyboard_layout (hashcat_ctx_t *hashcat_ctx, const char *filename, kb_layout_map_t *kb_layout_map, int *kb_layout_map_cnt) { char *line_buf = (char *) hcmalloc (HCBUFSIZ_LARGE); @@ -2641,8 +2649,7 @@ static bool initialize_keyboard_layout (hashcat_ctx_t *hashcat_ctx, const char * return false; } - u32 verifyF[256] = { 0 }; - u32 verifyT[256] = { 0 }; + int maps_cnt = 0; while (!feof (fp)) { @@ -2650,61 +2657,62 @@ static bool initialize_keyboard_layout (hashcat_ctx_t *hashcat_ctx, const char * if (line_len == 0) continue; - if (line_buf[1] != '=') + token_t token; + + token.token_cnt = 2; + + token.len_min[0] = 1; + token.len_max[0] = 4; + token.sep[0] = '='; + token.attr[0] = TOKEN_ATTR_VERIFY_LENGTH; + + token.len_min[1] = 0; + token.len_max[1] = 4; + token.sep[1] = '='; + token.attr[1] = TOKEN_ATTR_VERIFY_LENGTH; + + const int rc_tokenizer = input_tokenizer ((u8 *) line_buf, line_len, &token); + + if (rc_tokenizer != PARSER_OK) { event_log_error (hashcat_ctx, "%s: Syntax error: %s", filename, line_buf); + fclose (fp); + free (line_buf); return false; } - if (line_len == 2) - { - const u8 from = line_buf[0]; + memcpy (&kb_layout_map[maps_cnt].src_char, token.buf[0], token.len[0]); + memcpy (&kb_layout_map[maps_cnt].dst_char, token.buf[1], token.len[1]); - verifyF[from]++; - } - else if (line_len == 3) - { - const u8 from = line_buf[0]; - const u8 to = line_buf[2]; - - keyboard_layout[from] = to; + kb_layout_map[maps_cnt].src_len = token.len[0]; + kb_layout_map[maps_cnt].dst_len = token.len[1]; - verifyF[from]++; - verifyT[to]++; - } - else + if (maps_cnt == 256) { - event_log_error (hashcat_ctx, "%s: Syntax error: %s", filename, line_buf); + event_log_error (hashcat_ctx, "%s: too many entries", filename); + + fclose (fp); free (line_buf); return false; } + + maps_cnt++; } + *kb_layout_map_cnt = maps_cnt; + fclose (fp); free (line_buf); - for (int i = 0x20; i < 0x7f; i++) - { - if (verifyF[i] > 1) - { - event_log_error (hashcat_ctx, "%s: Mapping error: defined '%c' too often in from section", filename, i); - - return false; - } - - if (verifyT[i] > 1) - { - event_log_error (hashcat_ctx, "%s: Mapping error: defined '%c' too often in to section", filename, i); + // we need to sort this by length to ensure the largest blocks come first in mapping - return false; - } - } + qsort (kb_layout_map, maps_cnt, sizeof (kb_layout_map_t), sort_by_src_len); return true; } @@ -7026,11 +7034,6 @@ int truecrypt_parse_hash_1k (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAY if (entropy < MIN_SUFFICIENT_ENTROPY_FILE) return (PARSER_INSUFFICIENT_ENTROPY); - for (int i = 0; i < 256; i++) - { - tc->keyboard_layout[i] = i; - } - memcpy (tc->salt_buf, buf, 64); memcpy (tc->data_buf, buf + 64, 512 - 64); @@ -7074,11 +7077,6 @@ int truecrypt_parse_hash_2k (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAY if (entropy < MIN_SUFFICIENT_ENTROPY_FILE) return (PARSER_INSUFFICIENT_ENTROPY); - for (int i = 0; i < 256; i++) - { - tc->keyboard_layout[i] = i; - } - memcpy (tc->salt_buf, buf, 64); memcpy (tc->data_buf, buf + 64, 512 - 64); @@ -7122,11 +7120,6 @@ int veracrypt_parse_hash_200000 (u8 *input_buf, u32 input_len, hash_t *hash_buf, if (entropy < MIN_SUFFICIENT_ENTROPY_FILE) return (PARSER_INSUFFICIENT_ENTROPY); - for (int i = 0; i < 256; i++) - { - tc->keyboard_layout[i] = i; - } - memcpy (tc->salt_buf, buf, 64); memcpy (tc->data_buf, buf + 64, 512 - 64); @@ -7170,11 +7163,6 @@ int veracrypt_parse_hash_500000 (u8 *input_buf, u32 input_len, hash_t *hash_buf, if (entropy < MIN_SUFFICIENT_ENTROPY_FILE) return (PARSER_INSUFFICIENT_ENTROPY); - for (int i = 0; i < 256; i++) - { - tc->keyboard_layout[i] = i; - } - memcpy (tc->salt_buf, buf, 64); memcpy (tc->data_buf, buf + 64, 512 - 64); @@ -7218,11 +7206,6 @@ int veracrypt_parse_hash_327661 (u8 *input_buf, u32 input_len, hash_t *hash_buf, if (entropy < MIN_SUFFICIENT_ENTROPY_FILE) return (PARSER_INSUFFICIENT_ENTROPY); - for (int i = 0; i < 256; i++) - { - tc->keyboard_layout[i] = i; - } - memcpy (tc->salt_buf, buf, 64); memcpy (tc->data_buf, buf + 64, 512 - 64); @@ -7266,11 +7249,6 @@ int veracrypt_parse_hash_655331 (u8 *input_buf, u32 input_len, hash_t *hash_buf, if (entropy < MIN_SUFFICIENT_ENTROPY_FILE) return (PARSER_INSUFFICIENT_ENTROPY); - for (int i = 0; i < 256; i++) - { - tc->keyboard_layout[i] = i; - } - memcpy (tc->salt_buf, buf, 64); memcpy (tc->data_buf, buf + 64, 512 - 64); @@ -29327,7 +29305,7 @@ int hashconfig_general_defaults (hashcat_ctx_t *hashcat_ctx) if (optional_param2) { - const bool rc = initialize_keyboard_layout (hashcat_ctx, optional_param2, tc->keyboard_layout); + const bool rc = initialize_keyboard_layout (hashcat_ctx, optional_param2, tc->kb_layout_map, &tc->kb_layout_map_cnt); if (rc == false) return -1; }