mirror of
https://github.com/hashcat/hashcat.git
synced 2024-11-30 03:48:17 +00:00
Workaround added for NVidia OpenCL runtime: RACF kernel requires EBCDIC lookup to be done on shared memory
This commit is contained in:
jsteube
parent
2dd8018915
commit
c094f3b511
@ -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];
|
||||
|
||||
@ -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];
|
||||
|
||||
@ -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)
|
||||
|
||||
@ -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
|
||||
|
||||
|
||||
Loading…
Reference in New Issue
Block a user