arno
/
hashcat
mirror of https://github.com/hashcat/hashcat.git
1
0
Fork 0
Code Issues Releases Wiki Activity

Optimized -m 8500 for ROCm

Browse Source
pull/1309/head
jsteube 7 years ago
parent 5bcda7d05a
commit 02e2279d59
3 changed files with 40 additions and 82 deletions
Show Stats Download Patch File Download Diff File

38
OpenCL/m08500_a0-optimized.cl
Unescape View File

@ -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], __local u32 *s_ascii_to_ebcdic_pc)
void transform_racf_key (const u32x w0, const u32x w1, u32x key[2])
{
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), 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;
key[0] = BOX1 (((w0 >> 0) & 0xff), c_ascii_to_ebcdic_pc) << 0
| BOX1 (((w0 >> 8) & 0xff), c_ascii_to_ebcdic_pc) << 8
| BOX1 (((w0 >> 16) & 0xff), c_ascii_to_ebcdic_pc) << 16
| BOX1 (((w0 >> 24) & 0xff), c_ascii_to_ebcdic_pc) << 24;
key[1] = BOX1 (((w1 >> 0) & 0xff), c_ascii_to_ebcdic_pc) << 0
| BOX1 (((w1 >> 8) & 0xff), c_ascii_to_ebcdic_pc) << 8
| BOX1 (((w1 >> 16) & 0xff), c_ascii_to_ebcdic_pc) << 16
| BOX1 (((w1 >> 24) & 0xff), c_ascii_to_ebcdic_pc) << 24;
}
__kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_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,7 +535,6 @@ __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];
@ -560,13 +559,6 @@ __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;
/**
@ -615,7 +607,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, s_ascii_to_ebcdic_pc);
transform_racf_key (w0[0], w0[1], key);
const u32x c = key[0];
const u32x d = key[1];
@ -662,7 +654,6 @@ __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];
@ -687,11 +678,6 @@ __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;
@ -754,7 +740,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, s_ascii_to_ebcdic_pc);
transform_racf_key (w0[0], w0[1], key);
const u32x c = key[0];
const u32x d = key[1];

40
OpenCL/m08500_a1-optimized.cl
Unescape View File

@ -506,17 +506,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], __local u32 *s_ascii_to_ebcdic_pc)
void transform_racf_key (const u32x w0, const u32x w1, u32x key[2])
{
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), 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;
key[0] = BOX1 (((w0 >> 0) & 0xff), c_ascii_to_ebcdic_pc) << 0
| BOX1 (((w0 >> 8) & 0xff), c_ascii_to_ebcdic_pc) << 8
| BOX1 (((w0 >> 16) & 0xff), c_ascii_to_ebcdic_pc) << 16
| BOX1 (((w0 >> 24) & 0xff), c_ascii_to_ebcdic_pc) << 24;
key[1] = BOX1 (((w1 >> 0) & 0xff), c_ascii_to_ebcdic_pc) << 0
| BOX1 (((w1 >> 8) & 0xff), c_ascii_to_ebcdic_pc) << 8
| BOX1 (((w1 >> 16) & 0xff), c_ascii_to_ebcdic_pc) << 16
| BOX1 (((w1 >> 24) & 0xff), c_ascii_to_ebcdic_pc) << 24;
}
__kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_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)
@ -533,7 +533,6 @@ __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];
@ -558,13 +557,6 @@ __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;
/**
@ -656,7 +648,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, s_ascii_to_ebcdic_pc);
transform_racf_key (w0[0], w0[1], key);
const u32x c = key[0];
const u32x d = key[1];
@ -703,7 +695,6 @@ __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];
@ -728,13 +719,6 @@ __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;
/**
@ -838,7 +822,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, s_ascii_to_ebcdic_pc);
transform_racf_key (w0[0], w0[1], key);
const u32x c = key[0];
const u32x d = key[1];

44
OpenCL/m08500_a3-optimized.cl
Unescape View File

@ -506,20 +506,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], __local u32 *s_ascii_to_ebcdic_pc)
void transform_racf_key (const u32x w0, const u32x w1, u32x key[2])
{
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), 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;
key[0] = BOX1 (((w0 >> 0) & 0xff), c_ascii_to_ebcdic_pc) << 0
| BOX1 (((w0 >> 8) & 0xff), c_ascii_to_ebcdic_pc) << 8
| BOX1 (((w0 >> 16) & 0xff), c_ascii_to_ebcdic_pc) << 16
| BOX1 (((w0 >> 24) & 0xff), c_ascii_to_ebcdic_pc) << 24;
key[1] = BOX1 (((w1 >> 0) & 0xff), c_ascii_to_ebcdic_pc) << 0
| BOX1 (((w1 >> 8) & 0xff), c_ascii_to_ebcdic_pc) << 8
| BOX1 (((w1 >> 16) & 0xff), c_ascii_to_ebcdic_pc) << 16
| BOX1 (((w1 >> 24) & 0xff), c_ascii_to_ebcdic_pc) << 24;
}
void m08500m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], __local u32 *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 pw_t *combs_buf, __constant const 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], u32 w[16], const u32 pw_len, __global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const 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
@ -557,7 +557,7 @@ void m08500m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], __local u3
u32x key[2];
transform_racf_key (w0, w1, key, s_ascii_to_ebcdic_pc);
transform_racf_key (w0, w1, key);
const u32x c = key[0];
const u32x d = key[1];
@ -582,7 +582,7 @@ void m08500m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], __local u3
}
}
void m08500s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], __local u32 *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 pw_t *combs_buf, __constant const 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], u32 w[16], const u32 pw_len, __global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const 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
@ -632,7 +632,7 @@ void m08500s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], __local u3
u32x key[2];
transform_racf_key (w0, w1, key, s_ascii_to_ebcdic_pc);
transform_racf_key (w0, w1, key);
const u32x c = key[0];
const u32x d = key[1];
@ -671,7 +671,6 @@ __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];
@ -696,11 +695,6 @@ __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;
@ -734,7 +728,7 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* main
*/
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);
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);
}
__kernel void m08500_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const 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)
@ -759,7 +753,6 @@ __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];
@ -784,11 +777,6 @@ __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;
@ -822,7 +810,7 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* main
*/
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);
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);
}
__kernel void m08500_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const 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)

Write Preview
Loading…
Cancel
Save