From 49fc7d45b70f0e1ef2204fdcd71afe91878dd113 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Thu, 15 Nov 2018 22:29:03 +0100 Subject: [PATCH] - Keymaps: Added hashcat keyboard mapping us.hckmap (can be used as template) - Keymaps: Added hashcat keyboard mapping de.hckmap --- OpenCL/inc_truecrypt_keyboard.cl | 22 ++++++++ OpenCL/inc_truecrypt_keyfile.cl | 24 ++++++++ OpenCL/m06211-pure.cl | 50 +---------------- OpenCL/m06212-pure.cl | 50 +---------------- OpenCL/m06213-pure.cl | 50 +---------------- OpenCL/m06221-pure.cl | 55 +++++++++--------- OpenCL/m06222-pure.cl | 55 +++++++++--------- OpenCL/m06223-pure.cl | 55 +++++++++--------- OpenCL/m06231-pure.cl | 49 +++++++--------- OpenCL/m06232-pure.cl | 49 +++++++--------- OpenCL/m06233-pure.cl | 49 +++++++--------- OpenCL/m13751-pure.cl | 51 ++++++++--------- OpenCL/m13752-pure.cl | 51 ++++++++--------- OpenCL/m13753-pure.cl | 51 ++++++++--------- OpenCL/m13771-pure.cl | 49 ++++++---------- OpenCL/m13772-pure.cl | 49 ++++++---------- OpenCL/m13773-pure.cl | 49 ++++++---------- docs/changes.txt | 2 + layouts/de.hckmap | 95 ++++++++++++++++++++++++++++++++ layouts/us.hckmap | 95 ++++++++++++++++++++++++++++++++ src/interface.c | 41 +++++++------- 21 files changed, 529 insertions(+), 512 deletions(-) create mode 100644 OpenCL/inc_truecrypt_keyboard.cl create mode 100644 OpenCL/inc_truecrypt_keyfile.cl create mode 100644 layouts/de.hckmap create mode 100644 layouts/us.hckmap diff --git a/OpenCL/inc_truecrypt_keyboard.cl b/OpenCL/inc_truecrypt_keyboard.cl new file mode 100644 index 000000000..8d6fb7880 --- /dev/null +++ b/OpenCL/inc_truecrypt_keyboard.cl @@ -0,0 +1,22 @@ +DECLSPEC void keyboard_map (u32 w[4], __local u32 *s_keyboard_layout) +{ + 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); +} diff --git a/OpenCL/inc_truecrypt_keyfile.cl b/OpenCL/inc_truecrypt_keyfile.cl new file mode 100644 index 000000000..803bbed9d --- /dev/null +++ b/OpenCL/inc_truecrypt_keyfile.cl @@ -0,0 +1,24 @@ +DECLSPEC u32 u8add (const u32 a, const u32 b) +{ + const u32 a1 = (a >> 0) & 0xff; + const u32 a2 = (a >> 8) & 0xff; + const u32 a3 = (a >> 16) & 0xff; + const u32 a4 = (a >> 24) & 0xff; + + const u32 b1 = (b >> 0) & 0xff; + const u32 b2 = (b >> 8) & 0xff; + const u32 b3 = (b >> 16) & 0xff; + const u32 b4 = (b >> 24) & 0xff; + + const u32 r1 = (a1 + b1) & 0xff; + const u32 r2 = (a2 + b2) & 0xff; + const u32 r3 = (a3 + b3) & 0xff; + const u32 r4 = (a4 + b4) & 0xff; + + const u32 r = r1 << 0 + | r2 << 8 + | r3 << 16 + | r4 << 24; + + return r; +} diff --git a/OpenCL/m06211-pure.cl b/OpenCL/m06211-pure.cl index 20b117d77..de15de824 100644 --- a/OpenCL/m06211-pure.cl +++ b/OpenCL/m06211-pure.cl @@ -17,57 +17,11 @@ #include "inc_cipher_twofish.cl" #include "inc_cipher_serpent.cl" +#include "inc_truecrypt_keyfile.cl" +#include "inc_truecrypt_keyboard.cl" #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -DECLSPEC void keyboard_map (u32 w[4], __local u32 *s_keyboard_layout) -{ - 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); -} - -DECLSPEC u32 u8add (const u32 a, const u32 b) -{ - const u32 a1 = (a >> 0) & 0xff; - const u32 a2 = (a >> 8) & 0xff; - const u32 a3 = (a >> 16) & 0xff; - const u32 a4 = (a >> 24) & 0xff; - - const u32 b1 = (b >> 0) & 0xff; - const u32 b2 = (b >> 8) & 0xff; - const u32 b3 = (b >> 16) & 0xff; - const u32 b4 = (b >> 24) & 0xff; - - const u32 r1 = (a1 + b1) & 0xff; - const u32 r2 = (a2 + b2) & 0xff; - const u32 r3 = (a3 + b3) & 0xff; - const u32 r4 = (a4 + b4) & 0xff; - - const u32 r = r1 << 0 - | r2 << 8 - | r3 << 16 - | r4 << 24; - - return r; -} - DECLSPEC void hmac_ripemd160_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m06212-pure.cl b/OpenCL/m06212-pure.cl index c76261a13..0c4fa8211 100644 --- a/OpenCL/m06212-pure.cl +++ b/OpenCL/m06212-pure.cl @@ -17,57 +17,11 @@ #include "inc_cipher_twofish.cl" #include "inc_cipher_serpent.cl" +#include "inc_truecrypt_keyfile.cl" +#include "inc_truecrypt_keyboard.cl" #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -DECLSPEC void keyboard_map (u32 w[4], __local u32 *s_keyboard_layout) -{ - 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); -} - -DECLSPEC u32 u8add (const u32 a, const u32 b) -{ - const u32 a1 = (a >> 0) & 0xff; - const u32 a2 = (a >> 8) & 0xff; - const u32 a3 = (a >> 16) & 0xff; - const u32 a4 = (a >> 24) & 0xff; - - const u32 b1 = (b >> 0) & 0xff; - const u32 b2 = (b >> 8) & 0xff; - const u32 b3 = (b >> 16) & 0xff; - const u32 b4 = (b >> 24) & 0xff; - - const u32 r1 = (a1 + b1) & 0xff; - const u32 r2 = (a2 + b2) & 0xff; - const u32 r3 = (a3 + b3) & 0xff; - const u32 r4 = (a4 + b4) & 0xff; - - const u32 r = r1 << 0 - | r2 << 8 - | r3 << 16 - | r4 << 24; - - return r; -} - DECLSPEC void hmac_ripemd160_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m06213-pure.cl b/OpenCL/m06213-pure.cl index c5d7b6bf9..29582479b 100644 --- a/OpenCL/m06213-pure.cl +++ b/OpenCL/m06213-pure.cl @@ -17,57 +17,11 @@ #include "inc_cipher_twofish.cl" #include "inc_cipher_serpent.cl" +#include "inc_truecrypt_keyfile.cl" +#include "inc_truecrypt_keyboard.cl" #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -DECLSPEC void keyboard_map (u32 w[4], __local u32 *s_keyboard_layout) -{ - 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); -} - -DECLSPEC u32 u8add (const u32 a, const u32 b) -{ - const u32 a1 = (a >> 0) & 0xff; - const u32 a2 = (a >> 8) & 0xff; - const u32 a3 = (a >> 16) & 0xff; - const u32 a4 = (a >> 24) & 0xff; - - const u32 b1 = (b >> 0) & 0xff; - const u32 b2 = (b >> 8) & 0xff; - const u32 b3 = (b >> 16) & 0xff; - const u32 b4 = (b >> 24) & 0xff; - - const u32 r1 = (a1 + b1) & 0xff; - const u32 r2 = (a2 + b2) & 0xff; - const u32 r3 = (a3 + b3) & 0xff; - const u32 r4 = (a4 + b4) & 0xff; - - const u32 r = r1 << 0 - | r2 << 8 - | r3 << 16 - | r4 << 24; - - return r; -} - DECLSPEC void hmac_ripemd160_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m06221-pure.cl b/OpenCL/m06221-pure.cl index 790b35217..309771ec4 100644 --- a/OpenCL/m06221-pure.cl +++ b/OpenCL/m06221-pure.cl @@ -17,34 +17,11 @@ #include "inc_cipher_twofish.cl" #include "inc_cipher_serpent.cl" +#include "inc_truecrypt_keyfile.cl" +#include "inc_truecrypt_keyboard.cl" #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -DECLSPEC u32 u8add (const u32 a, const u32 b) -{ - const u32 a1 = (a >> 0) & 0xff; - const u32 a2 = (a >> 8) & 0xff; - const u32 a3 = (a >> 16) & 0xff; - const u32 a4 = (a >> 24) & 0xff; - - const u32 b1 = (b >> 0) & 0xff; - const u32 b2 = (b >> 8) & 0xff; - const u32 b3 = (b >> 16) & 0xff; - const u32 b4 = (b >> 24) & 0xff; - - const u32 r1 = (a1 + b1) & 0xff; - const u32 r2 = (a2 + b2) & 0xff; - const u32 r3 = (a3 + b3) & 0xff; - const u32 r4 = (a4 + b4) & 0xff; - - const u32 r = r1 << 0 - | r2 << 8 - | r3 << 16 - | r4 << 24; - - return r; -} - DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w4, u32x *w5, u32x *w6, u32x *w7, u64x *ipad, u64x *opad, u64x *digest) { digest[0] = ipad[0]; @@ -105,14 +82,29 @@ DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w __kernel void m06221_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc64_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { + const u64 gid = get_global_id (0); + const u64 lid = get_local_id (0); + const u64 lsz = get_local_size (0); + /** - * base + * keyboard layout shared */ - const u64 gid = get_global_id (0); + __local u32 s_keyboard_layout[256]; + + for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) + { + s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; + /** + * base + */ + u32 w0[4]; u32 w1[4]; u32 w2[4]; @@ -155,6 +147,15 @@ __kernel void m06221_init (__global pw_t *pws, __global const kernel_rule_t *rul 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); + w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]); diff --git a/OpenCL/m06222-pure.cl b/OpenCL/m06222-pure.cl index 9219b6f5c..a275d870a 100644 --- a/OpenCL/m06222-pure.cl +++ b/OpenCL/m06222-pure.cl @@ -17,34 +17,11 @@ #include "inc_cipher_twofish.cl" #include "inc_cipher_serpent.cl" +#include "inc_truecrypt_keyfile.cl" +#include "inc_truecrypt_keyboard.cl" #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -DECLSPEC u32 u8add (const u32 a, const u32 b) -{ - const u32 a1 = (a >> 0) & 0xff; - const u32 a2 = (a >> 8) & 0xff; - const u32 a3 = (a >> 16) & 0xff; - const u32 a4 = (a >> 24) & 0xff; - - const u32 b1 = (b >> 0) & 0xff; - const u32 b2 = (b >> 8) & 0xff; - const u32 b3 = (b >> 16) & 0xff; - const u32 b4 = (b >> 24) & 0xff; - - const u32 r1 = (a1 + b1) & 0xff; - const u32 r2 = (a2 + b2) & 0xff; - const u32 r3 = (a3 + b3) & 0xff; - const u32 r4 = (a4 + b4) & 0xff; - - const u32 r = r1 << 0 - | r2 << 8 - | r3 << 16 - | r4 << 24; - - return r; -} - DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w4, u32x *w5, u32x *w6, u32x *w7, u64x *ipad, u64x *opad, u64x *digest) { digest[0] = ipad[0]; @@ -105,14 +82,29 @@ DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w __kernel void m06222_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc64_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { + const u64 gid = get_global_id (0); + const u64 lid = get_local_id (0); + const u64 lsz = get_local_size (0); + /** - * base + * keyboard layout shared */ - const u64 gid = get_global_id (0); + __local u32 s_keyboard_layout[256]; + + for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) + { + s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; + /** + * base + */ + u32 w0[4]; u32 w1[4]; u32 w2[4]; @@ -155,6 +147,15 @@ __kernel void m06222_init (__global pw_t *pws, __global const kernel_rule_t *rul 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); + w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]); diff --git a/OpenCL/m06223-pure.cl b/OpenCL/m06223-pure.cl index 30e3e70c9..e0f97e5ed 100644 --- a/OpenCL/m06223-pure.cl +++ b/OpenCL/m06223-pure.cl @@ -17,34 +17,11 @@ #include "inc_cipher_twofish.cl" #include "inc_cipher_serpent.cl" +#include "inc_truecrypt_keyfile.cl" +#include "inc_truecrypt_keyboard.cl" #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -DECLSPEC u32 u8add (const u32 a, const u32 b) -{ - const u32 a1 = (a >> 0) & 0xff; - const u32 a2 = (a >> 8) & 0xff; - const u32 a3 = (a >> 16) & 0xff; - const u32 a4 = (a >> 24) & 0xff; - - const u32 b1 = (b >> 0) & 0xff; - const u32 b2 = (b >> 8) & 0xff; - const u32 b3 = (b >> 16) & 0xff; - const u32 b4 = (b >> 24) & 0xff; - - const u32 r1 = (a1 + b1) & 0xff; - const u32 r2 = (a2 + b2) & 0xff; - const u32 r3 = (a3 + b3) & 0xff; - const u32 r4 = (a4 + b4) & 0xff; - - const u32 r = r1 << 0 - | r2 << 8 - | r3 << 16 - | r4 << 24; - - return r; -} - DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w4, u32x *w5, u32x *w6, u32x *w7, u64x *ipad, u64x *opad, u64x *digest) { digest[0] = ipad[0]; @@ -105,14 +82,29 @@ DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w __kernel void m06223_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc64_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { + const u64 gid = get_global_id (0); + const u64 lid = get_local_id (0); + const u64 lsz = get_local_size (0); + /** - * base + * keyboard layout shared */ - const u64 gid = get_global_id (0); + __local u32 s_keyboard_layout[256]; + + for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) + { + s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; + /** + * base + */ + u32 w0[4]; u32 w1[4]; u32 w2[4]; @@ -155,6 +147,15 @@ __kernel void m06223_init (__global pw_t *pws, __global const kernel_rule_t *rul 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); + w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]); diff --git a/OpenCL/m06231-pure.cl b/OpenCL/m06231-pure.cl index 07804cda7..5ac3c0eb3 100644 --- a/OpenCL/m06231-pure.cl +++ b/OpenCL/m06231-pure.cl @@ -17,34 +17,11 @@ #include "inc_cipher_twofish.cl" #include "inc_cipher_serpent.cl" +#include "inc_truecrypt_keyfile.cl" +#include "inc_truecrypt_keyboard.cl" #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -DECLSPEC u32 u8add (const u32 a, const u32 b) -{ - const u32 a1 = (a >> 0) & 0xff; - const u32 a2 = (a >> 8) & 0xff; - const u32 a3 = (a >> 16) & 0xff; - const u32 a4 = (a >> 24) & 0xff; - - const u32 b1 = (b >> 0) & 0xff; - const u32 b2 = (b >> 8) & 0xff; - const u32 b3 = (b >> 16) & 0xff; - const u32 b4 = (b >> 24) & 0xff; - - const u32 r1 = (a1 + b1) & 0xff; - const u32 r2 = (a2 + b2) & 0xff; - const u32 r3 = (a3 + b3) & 0xff; - const u32 r4 = (a4 + b4) & 0xff; - - const u32 r = r1 << 0 - | r2 << 8 - | r3 << 16 - | r4 << 24; - - return r; -} - DECLSPEC void hmac_whirlpool_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest, __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256]) { digest[ 0] = ipad[ 0]; @@ -143,18 +120,21 @@ DECLSPEC void hmac_whirlpool_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x __kernel void m06231_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { - /** - * modifier - */ - const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); const u64 lsz = get_local_size (0); /** - * shared + * keyboard layout shared */ + __local u32 s_keyboard_layout[256]; + + for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) + { + s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + } + __local u32 s_Ch[8][256]; __local u32 s_Cl[8][256]; @@ -183,6 +163,10 @@ __kernel void m06231_init (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; + /** + * base + */ + u32 w0[4]; u32 w1[4]; u32 w2[4]; @@ -205,6 +189,11 @@ __kernel void m06231_init (__global pw_t *pws, __global const kernel_rule_t *rul 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); + w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]); diff --git a/OpenCL/m06232-pure.cl b/OpenCL/m06232-pure.cl index db1e591f2..4c015986c 100644 --- a/OpenCL/m06232-pure.cl +++ b/OpenCL/m06232-pure.cl @@ -17,34 +17,11 @@ #include "inc_cipher_twofish.cl" #include "inc_cipher_serpent.cl" +#include "inc_truecrypt_keyfile.cl" +#include "inc_truecrypt_keyboard.cl" #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -DECLSPEC u32 u8add (const u32 a, const u32 b) -{ - const u32 a1 = (a >> 0) & 0xff; - const u32 a2 = (a >> 8) & 0xff; - const u32 a3 = (a >> 16) & 0xff; - const u32 a4 = (a >> 24) & 0xff; - - const u32 b1 = (b >> 0) & 0xff; - const u32 b2 = (b >> 8) & 0xff; - const u32 b3 = (b >> 16) & 0xff; - const u32 b4 = (b >> 24) & 0xff; - - const u32 r1 = (a1 + b1) & 0xff; - const u32 r2 = (a2 + b2) & 0xff; - const u32 r3 = (a3 + b3) & 0xff; - const u32 r4 = (a4 + b4) & 0xff; - - const u32 r = r1 << 0 - | r2 << 8 - | r3 << 16 - | r4 << 24; - - return r; -} - DECLSPEC void hmac_whirlpool_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest, __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256]) { digest[ 0] = ipad[ 0]; @@ -143,18 +120,21 @@ DECLSPEC void hmac_whirlpool_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x __kernel void m06232_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { - /** - * modifier - */ - const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); const u64 lsz = get_local_size (0); /** - * shared + * keyboard layout shared */ + __local u32 s_keyboard_layout[256]; + + for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) + { + s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + } + __local u32 s_Ch[8][256]; __local u32 s_Cl[8][256]; @@ -183,6 +163,10 @@ __kernel void m06232_init (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; + /** + * base + */ + u32 w0[4]; u32 w1[4]; u32 w2[4]; @@ -205,6 +189,11 @@ __kernel void m06232_init (__global pw_t *pws, __global const kernel_rule_t *rul 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); + w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]); diff --git a/OpenCL/m06233-pure.cl b/OpenCL/m06233-pure.cl index 957cbf227..aa637e3ca 100644 --- a/OpenCL/m06233-pure.cl +++ b/OpenCL/m06233-pure.cl @@ -17,34 +17,11 @@ #include "inc_cipher_twofish.cl" #include "inc_cipher_serpent.cl" +#include "inc_truecrypt_keyfile.cl" +#include "inc_truecrypt_keyboard.cl" #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -DECLSPEC u32 u8add (const u32 a, const u32 b) -{ - const u32 a1 = (a >> 0) & 0xff; - const u32 a2 = (a >> 8) & 0xff; - const u32 a3 = (a >> 16) & 0xff; - const u32 a4 = (a >> 24) & 0xff; - - const u32 b1 = (b >> 0) & 0xff; - const u32 b2 = (b >> 8) & 0xff; - const u32 b3 = (b >> 16) & 0xff; - const u32 b4 = (b >> 24) & 0xff; - - const u32 r1 = (a1 + b1) & 0xff; - const u32 r2 = (a2 + b2) & 0xff; - const u32 r3 = (a3 + b3) & 0xff; - const u32 r4 = (a4 + b4) & 0xff; - - const u32 r = r1 << 0 - | r2 << 8 - | r3 << 16 - | r4 << 24; - - return r; -} - DECLSPEC void hmac_whirlpool_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest, __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256]) { digest[ 0] = ipad[ 0]; @@ -143,18 +120,21 @@ DECLSPEC void hmac_whirlpool_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x __kernel void m06233_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { - /** - * modifier - */ - const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); const u64 lsz = get_local_size (0); /** - * shared + * keyboard layout shared */ + __local u32 s_keyboard_layout[256]; + + for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) + { + s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + } + __local u32 s_Ch[8][256]; __local u32 s_Cl[8][256]; @@ -183,6 +163,10 @@ __kernel void m06233_init (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; + /** + * base + */ + u32 w0[4]; u32 w1[4]; u32 w2[4]; @@ -205,6 +189,11 @@ __kernel void m06233_init (__global pw_t *pws, __global const kernel_rule_t *rul 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); + w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]); diff --git a/OpenCL/m13751-pure.cl b/OpenCL/m13751-pure.cl index 3749eaf2d..b99348e11 100644 --- a/OpenCL/m13751-pure.cl +++ b/OpenCL/m13751-pure.cl @@ -17,34 +17,11 @@ #include "inc_cipher_twofish.cl" #include "inc_cipher_serpent.cl" +#include "inc_truecrypt_keyfile.cl" +#include "inc_truecrypt_keyboard.cl" #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -DECLSPEC u32 u8add (const u32 a, const u32 b) -{ - const u32 a1 = (a >> 0) & 0xff; - const u32 a2 = (a >> 8) & 0xff; - const u32 a3 = (a >> 16) & 0xff; - const u32 a4 = (a >> 24) & 0xff; - - const u32 b1 = (b >> 0) & 0xff; - const u32 b2 = (b >> 8) & 0xff; - const u32 b3 = (b >> 16) & 0xff; - const u32 b4 = (b >> 24) & 0xff; - - const u32 r1 = (a1 + b1) & 0xff; - const u32 r2 = (a2 + b2) & 0xff; - const u32 r3 = (a3 + b3) & 0xff; - const u32 r4 = (a4 + b4) & 0xff; - - const u32 r = r1 << 0 - | r2 << 8 - | r3 << 16 - | r4 << 24; - - return r; -} - DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; @@ -89,14 +66,29 @@ DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *i __kernel void m13751_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { + const u64 gid = get_global_id (0); + const u64 lid = get_local_id (0); + const u64 lsz = get_local_size (0); + /** - * base + * keyboard layout shared */ - const u64 gid = get_global_id (0); + __local u32 s_keyboard_layout[256]; + + for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) + { + s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; + /** + * base + */ + u32 w0[4]; u32 w1[4]; u32 w2[4]; @@ -119,6 +111,11 @@ __kernel void m13751_init (__global pw_t *pws, __global const kernel_rule_t *rul 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); + w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]); diff --git a/OpenCL/m13752-pure.cl b/OpenCL/m13752-pure.cl index 8edd845c0..14169b760 100644 --- a/OpenCL/m13752-pure.cl +++ b/OpenCL/m13752-pure.cl @@ -17,34 +17,11 @@ #include "inc_cipher_twofish.cl" #include "inc_cipher_serpent.cl" +#include "inc_truecrypt_keyfile.cl" +#include "inc_truecrypt_keyboard.cl" #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -DECLSPEC u32 u8add (const u32 a, const u32 b) -{ - const u32 a1 = (a >> 0) & 0xff; - const u32 a2 = (a >> 8) & 0xff; - const u32 a3 = (a >> 16) & 0xff; - const u32 a4 = (a >> 24) & 0xff; - - const u32 b1 = (b >> 0) & 0xff; - const u32 b2 = (b >> 8) & 0xff; - const u32 b3 = (b >> 16) & 0xff; - const u32 b4 = (b >> 24) & 0xff; - - const u32 r1 = (a1 + b1) & 0xff; - const u32 r2 = (a2 + b2) & 0xff; - const u32 r3 = (a3 + b3) & 0xff; - const u32 r4 = (a4 + b4) & 0xff; - - const u32 r = r1 << 0 - | r2 << 8 - | r3 << 16 - | r4 << 24; - - return r; -} - DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; @@ -89,14 +66,29 @@ DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *i __kernel void m13752_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { + const u64 gid = get_global_id (0); + const u64 lid = get_local_id (0); + const u64 lsz = get_local_size (0); + /** - * base + * keyboard layout shared */ - const u64 gid = get_global_id (0); + __local u32 s_keyboard_layout[256]; + + for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) + { + s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; + /** + * base + */ + u32 w0[4]; u32 w1[4]; u32 w2[4]; @@ -119,6 +111,11 @@ __kernel void m13752_init (__global pw_t *pws, __global const kernel_rule_t *rul 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); + w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]); diff --git a/OpenCL/m13753-pure.cl b/OpenCL/m13753-pure.cl index e0d3fd3c7..d2e6f1f74 100644 --- a/OpenCL/m13753-pure.cl +++ b/OpenCL/m13753-pure.cl @@ -17,34 +17,11 @@ #include "inc_cipher_twofish.cl" #include "inc_cipher_serpent.cl" +#include "inc_truecrypt_keyfile.cl" +#include "inc_truecrypt_keyboard.cl" #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -DECLSPEC u32 u8add (const u32 a, const u32 b) -{ - const u32 a1 = (a >> 0) & 0xff; - const u32 a2 = (a >> 8) & 0xff; - const u32 a3 = (a >> 16) & 0xff; - const u32 a4 = (a >> 24) & 0xff; - - const u32 b1 = (b >> 0) & 0xff; - const u32 b2 = (b >> 8) & 0xff; - const u32 b3 = (b >> 16) & 0xff; - const u32 b4 = (b >> 24) & 0xff; - - const u32 r1 = (a1 + b1) & 0xff; - const u32 r2 = (a2 + b2) & 0xff; - const u32 r3 = (a3 + b3) & 0xff; - const u32 r4 = (a4 + b4) & 0xff; - - const u32 r = r1 << 0 - | r2 << 8 - | r3 << 16 - | r4 << 24; - - return r; -} - DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; @@ -89,14 +66,29 @@ DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *i __kernel void m13753_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { + const u64 gid = get_global_id (0); + const u64 lid = get_local_id (0); + const u64 lsz = get_local_size (0); + /** - * base + * keyboard layout shared */ - const u64 gid = get_global_id (0); + __local u32 s_keyboard_layout[256]; + + for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) + { + s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; + /** + * base + */ + u32 w0[4]; u32 w1[4]; u32 w2[4]; @@ -119,6 +111,11 @@ __kernel void m13753_init (__global pw_t *pws, __global const kernel_rule_t *rul 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); + w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]); diff --git a/OpenCL/m13771-pure.cl b/OpenCL/m13771-pure.cl index fab249ef1..8377f2810 100644 --- a/OpenCL/m13771-pure.cl +++ b/OpenCL/m13771-pure.cl @@ -17,34 +17,11 @@ #include "inc_cipher_twofish.cl" #include "inc_cipher_serpent.cl" +#include "inc_truecrypt_keyfile.cl" +#include "inc_truecrypt_keyboard.cl" #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -DECLSPEC u32 u8add (const u32 a, const u32 b) -{ - const u32 a1 = (a >> 0) & 0xff; - const u32 a2 = (a >> 8) & 0xff; - const u32 a3 = (a >> 16) & 0xff; - const u32 a4 = (a >> 24) & 0xff; - - const u32 b1 = (b >> 0) & 0xff; - const u32 b2 = (b >> 8) & 0xff; - const u32 b3 = (b >> 16) & 0xff; - const u32 b4 = (b >> 24) & 0xff; - - const u32 r1 = (a1 + b1) & 0xff; - const u32 r2 = (a2 + b2) & 0xff; - const u32 r3 = (a3 + b3) & 0xff; - const u32 r4 = (a4 + b4) & 0xff; - - const u32 r = r1 << 0 - | r2 << 8 - | r3 << 16 - | r4 << 24; - - return r; -} - DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u64x *ipad_hash, u64x *opad_hash, u64x *ipad_raw, u64x *opad_raw, u64x *digest, SHM_TYPE u64a (*s_sbob_sl64)[256]) { const u64x nullbuf[8] = { 0 }; @@ -132,17 +109,16 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6 __kernel void m13771_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global vc64_sbog_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { - /** - * base - */ - const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); const u64 lsz = get_local_size (0); - /** - * shared lookup table - */ + __local u32 s_keyboard_layout[256]; + + for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) + { + s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + } #ifdef REAL_SHM @@ -170,6 +146,10 @@ __kernel void m13771_init (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; + /** + * base + */ + u32 w0[4]; u32 w1[4]; u32 w2[4]; @@ -192,6 +172,11 @@ __kernel void m13771_init (__global pw_t *pws, __global const kernel_rule_t *rul 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); + w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]); diff --git a/OpenCL/m13772-pure.cl b/OpenCL/m13772-pure.cl index d60193b64..e83a039ed 100644 --- a/OpenCL/m13772-pure.cl +++ b/OpenCL/m13772-pure.cl @@ -17,34 +17,11 @@ #include "inc_cipher_twofish.cl" #include "inc_cipher_serpent.cl" +#include "inc_truecrypt_keyfile.cl" +#include "inc_truecrypt_keyboard.cl" #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -DECLSPEC u32 u8add (const u32 a, const u32 b) -{ - const u32 a1 = (a >> 0) & 0xff; - const u32 a2 = (a >> 8) & 0xff; - const u32 a3 = (a >> 16) & 0xff; - const u32 a4 = (a >> 24) & 0xff; - - const u32 b1 = (b >> 0) & 0xff; - const u32 b2 = (b >> 8) & 0xff; - const u32 b3 = (b >> 16) & 0xff; - const u32 b4 = (b >> 24) & 0xff; - - const u32 r1 = (a1 + b1) & 0xff; - const u32 r2 = (a2 + b2) & 0xff; - const u32 r3 = (a3 + b3) & 0xff; - const u32 r4 = (a4 + b4) & 0xff; - - const u32 r = r1 << 0 - | r2 << 8 - | r3 << 16 - | r4 << 24; - - return r; -} - DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u64x *ipad_hash, u64x *opad_hash, u64x *ipad_raw, u64x *opad_raw, u64x *digest, SHM_TYPE u64a (*s_sbob_sl64)[256]) { const u64x nullbuf[8] = { 0 }; @@ -132,17 +109,16 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6 __kernel void m13772_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global vc64_sbog_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { - /** - * base - */ - const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); const u64 lsz = get_local_size (0); - /** - * shared lookup table - */ + __local u32 s_keyboard_layout[256]; + + for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) + { + s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + } #ifdef REAL_SHM @@ -170,6 +146,10 @@ __kernel void m13772_init (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; + /** + * base + */ + u32 w0[4]; u32 w1[4]; u32 w2[4]; @@ -192,6 +172,11 @@ __kernel void m13772_init (__global pw_t *pws, __global const kernel_rule_t *rul 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); + w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]); diff --git a/OpenCL/m13773-pure.cl b/OpenCL/m13773-pure.cl index 2e03c52a0..ab1f98bb1 100644 --- a/OpenCL/m13773-pure.cl +++ b/OpenCL/m13773-pure.cl @@ -17,34 +17,11 @@ #include "inc_cipher_twofish.cl" #include "inc_cipher_serpent.cl" +#include "inc_truecrypt_keyfile.cl" +#include "inc_truecrypt_keyboard.cl" #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -DECLSPEC u32 u8add (const u32 a, const u32 b) -{ - const u32 a1 = (a >> 0) & 0xff; - const u32 a2 = (a >> 8) & 0xff; - const u32 a3 = (a >> 16) & 0xff; - const u32 a4 = (a >> 24) & 0xff; - - const u32 b1 = (b >> 0) & 0xff; - const u32 b2 = (b >> 8) & 0xff; - const u32 b3 = (b >> 16) & 0xff; - const u32 b4 = (b >> 24) & 0xff; - - const u32 r1 = (a1 + b1) & 0xff; - const u32 r2 = (a2 + b2) & 0xff; - const u32 r3 = (a3 + b3) & 0xff; - const u32 r4 = (a4 + b4) & 0xff; - - const u32 r = r1 << 0 - | r2 << 8 - | r3 << 16 - | r4 << 24; - - return r; -} - DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u64x *ipad_hash, u64x *opad_hash, u64x *ipad_raw, u64x *opad_raw, u64x *digest, SHM_TYPE u64a (*s_sbob_sl64)[256]) { const u64x nullbuf[8] = { 0 }; @@ -132,17 +109,16 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6 __kernel void m13773_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global vc64_sbog_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { - /** - * base - */ - const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); const u64 lsz = get_local_size (0); - /** - * shared lookup table - */ + __local u32 s_keyboard_layout[256]; + + for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz) + { + s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i]; + } #ifdef REAL_SHM @@ -170,6 +146,10 @@ __kernel void m13773_init (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; + /** + * base + */ + u32 w0[4]; u32 w1[4]; u32 w2[4]; @@ -192,6 +172,11 @@ __kernel void m13773_init (__global pw_t *pws, __global const kernel_rule_t *rul 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); + w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]); w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]); w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]); diff --git a/docs/changes.txt b/docs/changes.txt index 76817a533..1b657f036 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -56,6 +56,8 @@ - Brain: Set --brain-client-features default from 3 to 2 - Dependencies: Added xxHash and OpenCL-Headers to deps/ in order to allow building hashcat from GitHub source release package - Dependencies: Removed gitmodules xxHash and OpenCL-Headers +- Keymaps: Added hashcat keyboard mapping us.hckmap (can be used as template) +- Keymaps: Added hashcat keyboard mapping de.hckmap - Hardware Monitor: Renamed --gpu-temp-abort to --hwmon-temp-abort - Hardware Monitor: Renamed --gpu-temp-disable to --hwmon-disable - Memory: Limit maximum host memory allocation depending on bitness diff --git a/layouts/de.hckmap b/layouts/de.hckmap new file mode 100644 index 000000000..5d13b4dd7 --- /dev/null +++ b/layouts/de.hckmap @@ -0,0 +1,95 @@ + = +!=! +"=@ +#=\ +$=$ +%=% +&=^ +'=| +(=* +)=( +*=} ++=] +,=, +-=/ +.=/ +/=& +0=0 +1=1 +2=2 +3=3 +4=4 +5=5 +6=6 +7=7 +8=8 +9=9 +:=> +;=< +<=\ +==) +>=| +?=_ +@= +A=A +B=B +C=C +D=D +E=E +F=F +G=G +H=H +I=I +J=J +K=K +L=L +M=M +N=N +O=O +P=P +Q=Q +R=R +S=S +T=T +U=U +V=V +W=W +X=X +Y=Z +Z=Y +[= +\= +]= +^=` +_=? +`=+ +a=a +b=b +c=c +d=d +e=e +f=f +g=g +h=h +i=i +j=j +k=k +l=l +m=m +n=n +o=o +p=p +q=q +r=r +s=s +t=t +u=u +v=v +w=w +x=x +y=z +z=y +{= +|= +}= +~= diff --git a/layouts/us.hckmap b/layouts/us.hckmap new file mode 100644 index 000000000..a66c54782 --- /dev/null +++ b/layouts/us.hckmap @@ -0,0 +1,95 @@ + = +!=! +"=" +#=# +$=$ +%=% +&=& +'=' +(=( +)=) +*=* ++=+ +,=, +-=- +.=. +/=/ +0=0 +1=1 +2=2 +3=3 +4=4 +5=5 +6=6 +7=7 +8=8 +9=9 +:=: +;=; +<=< +=== +>=> +?=? +@=@ +A=A +B=B +C=C +D=D +E=E +F=F +G=G +H=H +I=I +J=J +K=K +L=L +M=M +N=N +O=O +P=P +Q=Q +R=R +S=S +T=T +U=U +V=V +W=W +X=X +Y=Y +Z=Z +[=[ +\=\ +]=] +^=^ +_=_ +`=` +a=a +b=b +c=c +d=d +e=e +f=f +g=g +h=h +i=i +j=j +k=k +l=l +m=m +n=n +o=o +p=p +q=q +r=r +s=s +t=t +u=u +v=v +w=w +x=x +y=y +z=z +{={ +|=| +}=} +~=~ diff --git a/src/interface.c b/src/interface.c index 8f8802245..2eddcf892 100644 --- a/src/interface.c +++ b/src/interface.c @@ -2650,23 +2650,32 @@ static bool initialize_keyboard_layout (hashcat_ctx_t *hashcat_ctx, const char * if (line_len == 0) continue; - token_t token; + if (line_buf[1] != '=') + { + event_log_error (hashcat_ctx, "%s: Syntax error: %s", filename, line_buf); - token.token_cnt = 2; + free (line_buf); - token.len_min[0] = 1; - token.len_max[0] = 1; - token.sep[0] = '='; - token.attr[0] = TOKEN_ATTR_VERIFY_LENGTH; + return false; + } + + if (line_len == 2) + { + const u8 from = line_buf[0]; - token.len_min[1] = 1; - token.len_max[1] = 1; - token.sep[1] = '='; - token.attr[1] = TOKEN_ATTR_VERIFY_LENGTH; + verifyF[from]++; + } + else if (line_len == 3) + { + const u8 from = line_buf[0]; + const u8 to = line_buf[2]; - const int rc_tokenizer = input_tokenizer ((u8 *) line_buf, line_len, &token); + keyboard_layout[from] = to; - if (rc_tokenizer != PARSER_OK) + verifyF[from]++; + verifyT[to]++; + } + else { event_log_error (hashcat_ctx, "%s: Syntax error: %s", filename, line_buf); @@ -2674,14 +2683,6 @@ static bool initialize_keyboard_layout (hashcat_ctx_t *hashcat_ctx, const char * return false; } - - const u8 from = token.buf[0][0]; - const u8 to = token.buf[1][0]; - - keyboard_layout[from] = to; - - verifyF[from]++; - verifyT[to]++; } fclose (fp);