From c094f3b51111d2402ea5be307f75c9c8033960d7 Mon Sep 17 00:00:00 2001 From: jsteube Date: Sat, 25 Feb 2017 17:36:29 +0100 Subject: [PATCH] Workaround added for NVidia OpenCL runtime: RACF kernel requires EBCDIC lookup to be done on shared memory --- OpenCL/m08500_a0.cl | 36 ++++++++++++++++++++++++------------ OpenCL/m08500_a1.cl | 36 ++++++++++++++++++++++++------------ OpenCL/m08500_a3.cl | 44 ++++++++++++++++++++++++++++---------------- docs/changes.txt | 1 + 4 files changed, 77 insertions(+), 40 deletions(-) diff --git a/OpenCL/m08500_a0.cl b/OpenCL/m08500_a0.cl index d66ab1db5..77ccb3dcc 100644 --- a/OpenCL/m08500_a0.cl +++ b/OpenCL/m08500_a0.cl @@ -54,7 +54,7 @@ PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \ } -__constant u8 ascii_to_ebcdic_pc[256] = +__constant u32 c_ascii_to_ebcdic_pc[256] = { // little hack, can't crack 0-bytes in password, but who cares // 0xab, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5, @@ -510,17 +510,17 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 } } -void transform_racf_key (const u32x w0, const u32x w1, u32x key[2]) +void transform_racf_key (const u32x w0, const u32x w1, u32x key[2], __local u32 *s_ascii_to_ebcdic_pc) { - key[0] = BOX1 (((w0 >> 0) & 0xff), ascii_to_ebcdic_pc) << 0 - | BOX1 (((w0 >> 8) & 0xff), ascii_to_ebcdic_pc) << 8 - | BOX1 (((w0 >> 16) & 0xff), ascii_to_ebcdic_pc) << 16 - | BOX1 (((w0 >> 24) & 0xff), ascii_to_ebcdic_pc) << 24; + key[0] = BOX1 (((w0 >> 0) & 0xff), s_ascii_to_ebcdic_pc) << 0 + | BOX1 (((w0 >> 8) & 0xff), s_ascii_to_ebcdic_pc) << 8 + | BOX1 (((w0 >> 16) & 0xff), s_ascii_to_ebcdic_pc) << 16 + | BOX1 (((w0 >> 24) & 0xff), s_ascii_to_ebcdic_pc) << 24; - key[1] = BOX1 (((w1 >> 0) & 0xff), ascii_to_ebcdic_pc) << 0 - | BOX1 (((w1 >> 8) & 0xff), ascii_to_ebcdic_pc) << 8 - | BOX1 (((w1 >> 16) & 0xff), ascii_to_ebcdic_pc) << 16 - | BOX1 (((w1 >> 24) & 0xff), ascii_to_ebcdic_pc) << 24; + key[1] = BOX1 (((w1 >> 0) & 0xff), s_ascii_to_ebcdic_pc) << 0 + | BOX1 (((w1 >> 8) & 0xff), s_ascii_to_ebcdic_pc) << 8 + | BOX1 (((w1 >> 16) & 0xff), s_ascii_to_ebcdic_pc) << 16 + | BOX1 (((w1 >> 24) & 0xff), s_ascii_to_ebcdic_pc) << 24; } __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *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 void *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 u32 gid_max) @@ -537,6 +537,7 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule * shared */ + __local u32 s_ascii_to_ebcdic_pc[256]; __local u32 s_SPtrans[8][64]; __local u32 s_skb[8][64]; @@ -561,6 +562,11 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule s_skb[7][i] = c_skb[7][i]; } + for (u32 i = lid; i < 256; i += lsz) + { + s_ascii_to_ebcdic_pc[i] = c_ascii_to_ebcdic_pc[i]; + } + barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; @@ -611,7 +617,7 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule u32x key[2]; - transform_racf_key (w0[0], w0[1], key); + transform_racf_key (w0[0], w0[1], key, s_ascii_to_ebcdic_pc); const u32x c = key[0]; const u32x d = key[1]; @@ -658,6 +664,7 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule * shared */ + __local u32 s_ascii_to_ebcdic_pc[256]; __local u32 s_SPtrans[8][64]; __local u32 s_skb[8][64]; @@ -682,6 +689,11 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule s_skb[7][i] = c_skb[7][i]; } + for (u32 i = lid; i < 256; i += lsz) + { + s_ascii_to_ebcdic_pc[i] = c_ascii_to_ebcdic_pc[i]; + } + barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; @@ -744,7 +756,7 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule u32x key[2]; - transform_racf_key (w0[0], w0[1], key); + transform_racf_key (w0[0], w0[1], key, s_ascii_to_ebcdic_pc); const u32x c = key[0]; const u32x d = key[1]; diff --git a/OpenCL/m08500_a1.cl b/OpenCL/m08500_a1.cl index 292536a17..6ed49c241 100644 --- a/OpenCL/m08500_a1.cl +++ b/OpenCL/m08500_a1.cl @@ -52,7 +52,7 @@ PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \ } -__constant u8 ascii_to_ebcdic_pc[256] = +__constant u32 c_ascii_to_ebcdic_pc[256] = { // little hack, can't crack 0-bytes in password, but who cares // 0xab, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5, @@ -508,17 +508,17 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 } } -void transform_racf_key (const u32x w0, const u32x w1, u32x key[2]) +void transform_racf_key (const u32x w0, const u32x w1, u32x key[2], __local u32 *s_ascii_to_ebcdic_pc) { - key[0] = BOX1 (((w0 >> 0) & 0xff), ascii_to_ebcdic_pc) << 0 - | BOX1 (((w0 >> 8) & 0xff), ascii_to_ebcdic_pc) << 8 - | BOX1 (((w0 >> 16) & 0xff), ascii_to_ebcdic_pc) << 16 - | BOX1 (((w0 >> 24) & 0xff), ascii_to_ebcdic_pc) << 24; + key[0] = BOX1 (((w0 >> 0) & 0xff), s_ascii_to_ebcdic_pc) << 0 + | BOX1 (((w0 >> 8) & 0xff), s_ascii_to_ebcdic_pc) << 8 + | BOX1 (((w0 >> 16) & 0xff), s_ascii_to_ebcdic_pc) << 16 + | BOX1 (((w0 >> 24) & 0xff), s_ascii_to_ebcdic_pc) << 24; - key[1] = BOX1 (((w1 >> 0) & 0xff), ascii_to_ebcdic_pc) << 0 - | BOX1 (((w1 >> 8) & 0xff), ascii_to_ebcdic_pc) << 8 - | BOX1 (((w1 >> 16) & 0xff), ascii_to_ebcdic_pc) << 16 - | BOX1 (((w1 >> 24) & 0xff), ascii_to_ebcdic_pc) << 24; + key[1] = BOX1 (((w1 >> 0) & 0xff), s_ascii_to_ebcdic_pc) << 0 + | BOX1 (((w1 >> 8) & 0xff), s_ascii_to_ebcdic_pc) << 8 + | BOX1 (((w1 >> 16) & 0xff), s_ascii_to_ebcdic_pc) << 16 + | BOX1 (((w1 >> 24) & 0xff), s_ascii_to_ebcdic_pc) << 24; } __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *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 void *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 u32 gid_max) @@ -535,6 +535,7 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule * shared */ + __local u32 s_ascii_to_ebcdic_pc[256]; __local u32 s_SPtrans[8][64]; __local u32 s_skb[8][64]; @@ -559,6 +560,11 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule s_skb[7][i] = c_skb[7][i]; } + for (u32 i = lid; i < 256; i += lsz) + { + s_ascii_to_ebcdic_pc[i] = c_ascii_to_ebcdic_pc[i]; + } + barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; @@ -652,7 +658,7 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule u32x key[2]; - transform_racf_key (w0[0], w0[1], key); + transform_racf_key (w0[0], w0[1], key, s_ascii_to_ebcdic_pc); const u32x c = key[0]; const u32x d = key[1]; @@ -699,6 +705,7 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule * shared */ + __local u32 s_ascii_to_ebcdic_pc[256]; __local u32 s_SPtrans[8][64]; __local u32 s_skb[8][64]; @@ -723,6 +730,11 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule s_skb[7][i] = c_skb[7][i]; } + for (u32 i = lid; i < 256; i += lsz) + { + s_ascii_to_ebcdic_pc[i] = c_ascii_to_ebcdic_pc[i]; + } + barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; @@ -828,7 +840,7 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule u32x key[2]; - transform_racf_key (w0[0], w0[1], key); + transform_racf_key (w0[0], w0[1], key, s_ascii_to_ebcdic_pc); const u32x c = key[0]; const u32x d = key[1]; diff --git a/OpenCL/m08500_a3.cl b/OpenCL/m08500_a3.cl index 5b1bd35d8..b7b8c541a 100644 --- a/OpenCL/m08500_a3.cl +++ b/OpenCL/m08500_a3.cl @@ -52,7 +52,7 @@ PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \ } -__constant u8 ascii_to_ebcdic_pc[256] = +__constant u32 c_ascii_to_ebcdic_pc[256] = { // little hack, can't crack 0-bytes in password, but who cares // 0xab, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5, @@ -508,20 +508,20 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 } } -void transform_racf_key (const u32x w0, const u32x w1, u32x key[2]) +void transform_racf_key (const u32x w0, const u32x w1, u32x key[2], __local u32 *s_ascii_to_ebcdic_pc) { - key[0] = BOX1 (((w0 >> 0) & 0xff), ascii_to_ebcdic_pc) << 0 - | BOX1 (((w0 >> 8) & 0xff), ascii_to_ebcdic_pc) << 8 - | BOX1 (((w0 >> 16) & 0xff), ascii_to_ebcdic_pc) << 16 - | BOX1 (((w0 >> 24) & 0xff), ascii_to_ebcdic_pc) << 24; + key[0] = BOX1 (((w0 >> 0) & 0xff), s_ascii_to_ebcdic_pc) << 0 + | BOX1 (((w0 >> 8) & 0xff), s_ascii_to_ebcdic_pc) << 8 + | BOX1 (((w0 >> 16) & 0xff), s_ascii_to_ebcdic_pc) << 16 + | BOX1 (((w0 >> 24) & 0xff), s_ascii_to_ebcdic_pc) << 24; - key[1] = BOX1 (((w1 >> 0) & 0xff), ascii_to_ebcdic_pc) << 0 - | BOX1 (((w1 >> 8) & 0xff), ascii_to_ebcdic_pc) << 8 - | BOX1 (((w1 >> 16) & 0xff), ascii_to_ebcdic_pc) << 16 - | BOX1 (((w1 >> 24) & 0xff), ascii_to_ebcdic_pc) << 24; + key[1] = BOX1 (((w1 >> 0) & 0xff), s_ascii_to_ebcdic_pc) << 0 + | BOX1 (((w1 >> 8) & 0xff), s_ascii_to_ebcdic_pc) << 8 + | BOX1 (((w1 >> 16) & 0xff), s_ascii_to_ebcdic_pc) << 16 + | BOX1 (((w1 >> 24) & 0xff), s_ascii_to_ebcdic_pc) << 24; } -void m08500m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16], const u32 pw_len, __global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __constant u32x * words_buf_r, __global void *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 void *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) +void m08500m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], __local *s_ascii_to_ebcdic_pc, u32 w[16], const u32 pw_len, __global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __constant u32x * words_buf_r, __global void *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 void *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) { /** * modifier @@ -559,7 +559,7 @@ void m08500m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16], u32x key[2]; - transform_racf_key (w0, w1, key); + transform_racf_key (w0, w1, key, s_ascii_to_ebcdic_pc); const u32x c = key[0]; const u32x d = key[1]; @@ -584,7 +584,7 @@ void m08500m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16], } } -void m08500s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16], const u32 pw_len, __global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __constant u32x * words_buf_r, __global void *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 void *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) +void m08500s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], __local *s_ascii_to_ebcdic_pc, u32 w[16], const u32 pw_len, __global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __constant u32x * words_buf_r, __global void *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 void *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) { /** * modifier @@ -634,7 +634,7 @@ void m08500s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16], u32x key[2]; - transform_racf_key (w0, w1, key); + transform_racf_key (w0, w1, key, s_ascii_to_ebcdic_pc); const u32x c = key[0]; const u32x d = key[1]; @@ -673,6 +673,7 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule * shared */ + __local u32 s_ascii_to_ebcdic_pc[256]; __local u32 s_SPtrans[8][64]; __local u32 s_skb[8][64]; @@ -697,6 +698,11 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule s_skb[7][i] = c_skb[7][i]; } + for (u32 i = lid; i < 256; i += lsz) + { + s_ascii_to_ebcdic_pc[i] = c_ascii_to_ebcdic_pc[i]; + } + barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; @@ -730,7 +736,7 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule * main */ - m08500m (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); + m08500m (s_SPtrans, s_skb, s_ascii_to_ebcdic_pc, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); } __kernel void m08500_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __constant u32x * words_buf_r, __global void *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 void *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 u32 gid_max) @@ -755,6 +761,7 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule * shared */ + __local u32 s_ascii_to_ebcdic_pc[256]; __local u32 s_SPtrans[8][64]; __local u32 s_skb[8][64]; @@ -779,6 +786,11 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule s_skb[7][i] = c_skb[7][i]; } + for (u32 i = lid; i < 256; i += lsz) + { + s_ascii_to_ebcdic_pc[i] = c_ascii_to_ebcdic_pc[i]; + } + barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; @@ -812,7 +824,7 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule * main */ - m08500s (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); + m08500s (s_SPtrans, s_skb, s_ascii_to_ebcdic_pc, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); } __kernel void m08500_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __constant u32x * words_buf_r, __global void *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 void *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 u32 gid_max) diff --git a/docs/changes.txt b/docs/changes.txt index 92c413d7a..9d07b516e 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -32,6 +32,7 @@ ## Workarounds ## +- Workaround added for NVidia OpenCL runtime: RACF kernel requires EBCDIC lookup to be done on shared memory - Workaround added for AMDGPU-Pro OpenCL runtime: AES encrypt and decrypt Invertkey function was calculated wrong in certain cases - Workaround added for AMDGPU-Pro OpenCL runtime: RAR3 kernel require a volatile variable to work correctly