github
/
hashcat
mirror of https://github.com/hashcat/hashcat
1
Fork 0
Code Releases Activity

No longer need reqd_work_group_size() in OpenCL kernels since modules taken care of this using threads_min and threads_max

This commit is contained in:
jsteube 2019-03-10 12:13:14 +01:00
parent 9da14d55a2
commit 26d7602bbd
46 changed files with 197 additions and 279 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
OpenCL/m03200-pure.cl
View File

@ -361,7 +361,7 @@ DECLSPEC void expand_key (u32 *E, u32 *W, const int len)
}
}
__kernel void __attribute__((reqd_work_group_size(8, 1, 1))) m03200_init (KERN_ATTR_TMPS (bcrypt_tmp_t))
__kernel void m03200_init (KERN_ATTR_TMPS (bcrypt_tmp_t))
{
/**
* base
@ -580,7 +580,7 @@ __kernel void __attribute__((reqd_work_group_size(8, 1, 1))) m03200_init (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(8, 1, 1))) m03200_loop (KERN_ATTR_TMPS (bcrypt_tmp_t))
__kernel void m03200_loop (KERN_ATTR_TMPS (bcrypt_tmp_t))
{
/**
* base
@ -778,7 +778,7 @@ __kernel void __attribute__((reqd_work_group_size(8, 1, 1))) m03200_loop (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(8, 1, 1))) m03200_comp (KERN_ATTR_TMPS (bcrypt_tmp_t))
__kernel void m03200_comp (KERN_ATTR_TMPS (bcrypt_tmp_t))
{
/**
* base

12
OpenCL/m07500_a0-optimized.cl
View File

@ -394,7 +394,7 @@ DECLSPEC void kerb_prepare (const u32 *w0, const u32 *w1, const u32 pw_len, cons
hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_m04 (KERN_ATTR_RULES_ESALT (krb5pa_t))
__kernel void m07500_m04 (KERN_ATTR_RULES_ESALT (krb5pa_t))
{
/**
* modifier
@ -492,15 +492,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_m08 (KERN_ATTR_RULES_ESALT (krb5pa_t))
__kernel void m07500_m08 (KERN_ATTR_RULES_ESALT (krb5pa_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_m16 (KERN_ATTR_RULES_ESALT (krb5pa_t))
__kernel void m07500_m16 (KERN_ATTR_RULES_ESALT (krb5pa_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_s04 (KERN_ATTR_RULES_ESALT (krb5pa_t))
__kernel void m07500_s04 (KERN_ATTR_RULES_ESALT (krb5pa_t))
{
/**
* modifier
@ -598,10 +598,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_s08 (KERN_ATTR_RULES_ESALT (krb5pa_t))
__kernel void m07500_s08 (KERN_ATTR_RULES_ESALT (krb5pa_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_s16 (KERN_ATTR_RULES_ESALT (krb5pa_t))
__kernel void m07500_s16 (KERN_ATTR_RULES_ESALT (krb5pa_t))
{
}

4
OpenCL/m07500_a0-pure.cl
View File

@ -278,7 +278,7 @@ DECLSPEC void kerb_prepare (const u32 *K, const u32 *checksum, u32 *digest)
digest[3] = ctx.opad.h[3];
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_mxx (KERN_ATTR_RULES_ESALT (krb5pa_t))
__kernel void m07500_mxx (KERN_ATTR_RULES_ESALT (krb5pa_t))
{
/**
* modifier
@ -349,7 +349,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_mxx (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_sxx (KERN_ATTR_RULES_ESALT (krb5pa_t))
__kernel void m07500_sxx (KERN_ATTR_RULES_ESALT (krb5pa_t))
{
/**
* modifier

12
OpenCL/m07500_a1-optimized.cl
View File

@ -392,7 +392,7 @@ DECLSPEC void kerb_prepare (const u32 *w0, const u32 *w1, const u32 pw_len, cons
hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_m04 (KERN_ATTR_ESALT (krb5pa_t))
__kernel void m07500_m04 (KERN_ATTR_ESALT (krb5pa_t))
{
/**
* modifier
@ -540,15 +540,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_m08 (KERN_ATTR_ESALT (krb5pa_t))
__kernel void m07500_m08 (KERN_ATTR_ESALT (krb5pa_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_m16 (KERN_ATTR_ESALT (krb5pa_t))
__kernel void m07500_m16 (KERN_ATTR_ESALT (krb5pa_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_s04 (KERN_ATTR_ESALT (krb5pa_t))
__kernel void m07500_s04 (KERN_ATTR_ESALT (krb5pa_t))
{
/**
* modifier
@ -696,10 +696,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_s08 (KERN_ATTR_ESALT (krb5pa_t))
__kernel void m07500_s08 (KERN_ATTR_ESALT (krb5pa_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_s16 (KERN_ATTR_ESALT (krb5pa_t))
__kernel void m07500_s16 (KERN_ATTR_ESALT (krb5pa_t))
{
}

4
OpenCL/m07500_a1-pure.cl
View File

@ -276,7 +276,7 @@ DECLSPEC void kerb_prepare (const u32 *K, const u32 *checksum, u32 *digest)
digest[3] = ctx.opad.h[3];
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_mxx (KERN_ATTR_ESALT (krb5pa_t))
__kernel void m07500_mxx (KERN_ATTR_ESALT (krb5pa_t))
{
/**
* modifier
@ -345,7 +345,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_mxx (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_sxx (KERN_ATTR_ESALT (krb5pa_t))
__kernel void m07500_sxx (KERN_ATTR_ESALT (krb5pa_t))
{
/**
* modifier

12
OpenCL/m07500_a3-optimized.cl
View File

@ -482,7 +482,7 @@ DECLSPEC void m07500 (__local RC4_KEY *rc4_key, u32 *w0, u32 *w1, u32 *w2, u32 *
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_m04 (KERN_ATTR_ESALT (krb5pa_t))
__kernel void m07500_m04 (KERN_ATTR_ESALT (krb5pa_t))
{
/**
* base
@ -534,7 +534,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_m04 (KERN_A
m07500 (rc4_key, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_m08 (KERN_ATTR_ESALT (krb5pa_t))
__kernel void m07500_m08 (KERN_ATTR_ESALT (krb5pa_t))
{
/**
* base
@ -586,11 +586,11 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_m08 (KERN_A
m07500 (rc4_key, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_m16 (KERN_ATTR_ESALT (krb5pa_t))
__kernel void m07500_m16 (KERN_ATTR_ESALT (krb5pa_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_s04 (KERN_ATTR_ESALT (krb5pa_t))
__kernel void m07500_s04 (KERN_ATTR_ESALT (krb5pa_t))
{
/**
* base
@ -642,7 +642,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_s04 (KERN_A
m07500 (rc4_key, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_s08 (KERN_ATTR_ESALT (krb5pa_t))
__kernel void m07500_s08 (KERN_ATTR_ESALT (krb5pa_t))
{
/**
* base
@ -694,6 +694,6 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_s08 (KERN_A
m07500 (rc4_key, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_s16 (KERN_ATTR_ESALT (krb5pa_t))
__kernel void m07500_s16 (KERN_ATTR_ESALT (krb5pa_t))
{
}

4
OpenCL/m07500_a3-pure.cl
View File

@ -276,7 +276,7 @@ DECLSPEC void kerb_prepare (const u32 *K, const u32 *checksum, u32 *digest)
digest[3] = ctx.opad.h[3];
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_mxx (KERN_ATTR_VECTOR_ESALT (krb5pa_t))
__kernel void m07500_mxx (KERN_ATTR_VECTOR_ESALT (krb5pa_t))
{
/**
* modifier
@ -370,7 +370,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_mxx (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_sxx (KERN_ATTR_VECTOR_ESALT (krb5pa_t))
__kernel void m07500_sxx (KERN_ATTR_VECTOR_ESALT (krb5pa_t))
{
/**
* modifier

2
OpenCL/m08900-pure.cl
View File

@ -215,8 +215,6 @@ DECLSPEC void scrypt_smix (uint4 *X, uint4 *T, __global uint4 * restrict V0, __g
}
}
// there can be no __attribute__((reqd_work_group_size(16, 1, 1))) because kernel is used by both -m 8900 and -m 9300
__kernel void m08900_init (KERN_ATTR_TMPS (scrypt_tmp_t))
{
/**

6
OpenCL/m09000-pure.cl
View File

@ -512,7 +512,7 @@ DECLSPEC void sha1_transform (const u32 *w0, const u32 *w1, const u32 *w2, const
digest[4] += E;
}
__kernel void __attribute__((reqd_work_group_size(8, 1, 1))) m09000_init (KERN_ATTR_TMPS (pwsafe2_tmp_t))
__kernel void m09000_init (KERN_ATTR_TMPS (pwsafe2_tmp_t))
{
/**
* base
@ -731,7 +731,7 @@ __kernel void __attribute__((reqd_work_group_size(8, 1, 1))) m09000_init (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(8, 1, 1))) m09000_loop (KERN_ATTR_TMPS (pwsafe2_tmp_t))
__kernel void m09000_loop (KERN_ATTR_TMPS (pwsafe2_tmp_t))
{
/**
* base
@ -793,7 +793,7 @@ __kernel void __attribute__((reqd_work_group_size(8, 1, 1))) m09000_loop (KERN_A
tmps[gid].digest[1] = R0;
}
__kernel void __attribute__((reqd_work_group_size(8, 1, 1))) m09000_comp (KERN_ATTR_TMPS (pwsafe2_tmp_t))
__kernel void m09000_comp (KERN_ATTR_TMPS (pwsafe2_tmp_t))
{
/**
* base

12
OpenCL/m09700_a0-optimized.cl
View File

@ -492,7 +492,7 @@ DECLSPEC void gen336 (u32 *digest_pre, u32 *salt_buf, u32 *digest)
md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_m04 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
__kernel void m09700_m04 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
{
/**
* modifier
@ -673,15 +673,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_m08 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
__kernel void m09700_m08 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_m16 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
__kernel void m09700_m16 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_s04 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
__kernel void m09700_s04 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
{
/**
* modifier
@ -874,10 +874,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_s08 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
__kernel void m09700_s08 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_s16 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
__kernel void m09700_s16 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
{
}

12
OpenCL/m09700_a1-optimized.cl
View File

@ -490,7 +490,7 @@ DECLSPEC void gen336 (u32 *digest_pre, u32 *salt_buf, u32 *digest)
md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_m04 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09700_m04 (KERN_ATTR_ESALT (oldoffice01_t))
{
/**
* modifier
@ -729,15 +729,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_m08 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09700_m08 (KERN_ATTR_ESALT (oldoffice01_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_m16 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09700_m16 (KERN_ATTR_ESALT (oldoffice01_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_s04 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09700_s04 (KERN_ATTR_ESALT (oldoffice01_t))
{
/**
* modifier
@ -988,10 +988,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_s08 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09700_s08 (KERN_ATTR_ESALT (oldoffice01_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_s16 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09700_s16 (KERN_ATTR_ESALT (oldoffice01_t))
{
}

12
OpenCL/m09700_a3-optimized.cl
View File

@ -975,7 +975,7 @@ DECLSPEC void m09700s (__local RC4_KEY *rc4_keys, u32 *w0, u32 *w1, u32 *w2, u32
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_m04 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09700_m04 (KERN_ATTR_ESALT (oldoffice01_t))
{
/**
* base
@ -1024,7 +1024,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_m04 (KERN_A
m09700m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_m08 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09700_m08 (KERN_ATTR_ESALT (oldoffice01_t))
{
/**
* base
@ -1073,7 +1073,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_m08 (KERN_A
m09700m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_m16 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09700_m16 (KERN_ATTR_ESALT (oldoffice01_t))
{
/**
* base
@ -1122,7 +1122,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_m16 (KERN_A
m09700m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_s04 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09700_s04 (KERN_ATTR_ESALT (oldoffice01_t))
{
/**
* base
@ -1171,7 +1171,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_s04 (KERN_A
m09700s (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_s08 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09700_s08 (KERN_ATTR_ESALT (oldoffice01_t))
{
/**
* base
@ -1220,7 +1220,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_s08 (KERN_A
m09700s (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_s16 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09700_s16 (KERN_ATTR_ESALT (oldoffice01_t))
{
/**
* base

12
OpenCL/m09710_a0-optimized.cl
View File

@ -148,7 +148,7 @@ DECLSPEC u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, const u32 *in, u3
return j;
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_m04 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
__kernel void m09710_m04 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
{
/**
* modifier
@ -280,15 +280,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_m08 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
__kernel void m09710_m08 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_m16 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
__kernel void m09710_m16 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_s04 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
__kernel void m09710_s04 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
{
/**
* modifier
@ -432,10 +432,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_s08 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
__kernel void m09710_s08 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_s16 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
__kernel void m09710_s16 (KERN_ATTR_RULES_ESALT (oldoffice01_t))
{
}

12
OpenCL/m09710_a1-optimized.cl
View File

@ -146,7 +146,7 @@ DECLSPEC u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, const u32 *in, u3
return j;
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_m04 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09710_m04 (KERN_ATTR_ESALT (oldoffice01_t))
{
/**
* modifier
@ -324,15 +324,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_m08 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09710_m08 (KERN_ATTR_ESALT (oldoffice01_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_m16 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09710_m16 (KERN_ATTR_ESALT (oldoffice01_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_s04 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09710_s04 (KERN_ATTR_ESALT (oldoffice01_t))
{
/**
* modifier
@ -522,10 +522,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_s08 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09710_s08 (KERN_ATTR_ESALT (oldoffice01_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_s16 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09710_s16 (KERN_ATTR_ESALT (oldoffice01_t))
{
}

12
OpenCL/m09710_a3-optimized.cl
View File

@ -380,7 +380,7 @@ DECLSPEC void m09710s (__local RC4_KEY *rc4_keys, u32 *w0, u32 *w1, u32 *w2, u32
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_m04 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09710_m04 (KERN_ATTR_ESALT (oldoffice01_t))
{
/**
* base
@ -429,15 +429,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_m04 (KERN_A
m09710m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_m08 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09710_m08 (KERN_ATTR_ESALT (oldoffice01_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_m16 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09710_m16 (KERN_ATTR_ESALT (oldoffice01_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_s04 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09710_s04 (KERN_ATTR_ESALT (oldoffice01_t))
{
/**
* base
@ -486,10 +486,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_s04 (KERN_A
m09710s (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_s08 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09710_s08 (KERN_ATTR_ESALT (oldoffice01_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_s16 (KERN_ATTR_ESALT (oldoffice01_t))
__kernel void m09710_s16 (KERN_ATTR_ESALT (oldoffice01_t))
{
}

12
OpenCL/m09800_a0-optimized.cl
View File

@ -148,7 +148,7 @@ DECLSPEC u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, const u32 *in, u3
return j;
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_m04 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
__kernel void m09800_m04 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
{
/**
* modifier
@ -340,15 +340,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_m08 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
__kernel void m09800_m08 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_m16 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
__kernel void m09800_m16 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_s04 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
__kernel void m09800_s04 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
{
/**
* modifier
@ -552,10 +552,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_s08 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
__kernel void m09800_s08 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_s16 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
__kernel void m09800_s16 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
{
}

12
OpenCL/m09800_a1-optimized.cl
View File

@ -146,7 +146,7 @@ DECLSPEC u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, const u32 *in, u3
return j;
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_m04 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09800_m04 (KERN_ATTR_ESALT (oldoffice34_t))
{
/**
* modifier
@ -388,15 +388,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_m08 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09800_m08 (KERN_ATTR_ESALT (oldoffice34_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_m16 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09800_m16 (KERN_ATTR_ESALT (oldoffice34_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_s04 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09800_s04 (KERN_ATTR_ESALT (oldoffice34_t))
{
/**
* modifier
@ -650,10 +650,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_s08 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09800_s08 (KERN_ATTR_ESALT (oldoffice34_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_s16 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09800_s16 (KERN_ATTR_ESALT (oldoffice34_t))
{
}

12
OpenCL/m09800_a3-optimized.cl
View File

@ -491,7 +491,7 @@ DECLSPEC void m09800s (__local RC4_KEY *rc4_keys, u32 *w0, u32 *w1, u32 *w2, u32
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_m04 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09800_m04 (KERN_ATTR_ESALT (oldoffice34_t))
{
/**
* base
@ -540,7 +540,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_m04 (KERN_A
m09800m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_m08 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09800_m08 (KERN_ATTR_ESALT (oldoffice34_t))
{
/**
* base
@ -589,7 +589,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_m08 (KERN_A
m09800m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_m16 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09800_m16 (KERN_ATTR_ESALT (oldoffice34_t))
{
/**
* base
@ -638,7 +638,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_m16 (KERN_A
m09800m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_s04 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09800_s04 (KERN_ATTR_ESALT (oldoffice34_t))
{
/**
* base
@ -687,7 +687,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_s04 (KERN_A
m09800s (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_s08 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09800_s08 (KERN_ATTR_ESALT (oldoffice34_t))
{
/**
* base
@ -736,7 +736,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_s08 (KERN_A
m09800s (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_s16 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09800_s16 (KERN_ATTR_ESALT (oldoffice34_t))
{
/**
* base

12
OpenCL/m09810_a0-optimized.cl
View File

@ -148,7 +148,7 @@ DECLSPEC u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, const u32 *in, u3
return j;
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_m04 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
__kernel void m09810_m04 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
{
/**
* modifier
@ -267,15 +267,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_m08 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
__kernel void m09810_m08 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_m16 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
__kernel void m09810_m16 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_s04 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
__kernel void m09810_s04 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
{
/**
* modifier
@ -406,10 +406,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_s08 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
__kernel void m09810_s08 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_s16 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
__kernel void m09810_s16 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
{
}

12
OpenCL/m09810_a1-optimized.cl
View File

@ -146,7 +146,7 @@ DECLSPEC u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, const u32 *in, u3
return j;
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_m04 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09810_m04 (KERN_ATTR_ESALT (oldoffice34_t))
{
/**
* modifier
@ -311,15 +311,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_m08 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09810_m08 (KERN_ATTR_ESALT (oldoffice34_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_m16 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09810_m16 (KERN_ATTR_ESALT (oldoffice34_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_s04 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09810_s04 (KERN_ATTR_ESALT (oldoffice34_t))
{
/**
* modifier
@ -496,10 +496,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_s08 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09810_s08 (KERN_ATTR_ESALT (oldoffice34_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_s16 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09810_s16 (KERN_ATTR_ESALT (oldoffice34_t))
{
}

12
OpenCL/m09810_a3-optimized.cl
View File

@ -350,7 +350,7 @@ DECLSPEC void m09810s (__local RC4_KEY *rc4_keys, u32 *w0, u32 *w1, u32 *w2, u32
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_m04 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09810_m04 (KERN_ATTR_ESALT (oldoffice34_t))
{
/**
* base
@ -399,7 +399,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_m04 (KERN_A
m09810m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_m08 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09810_m08 (KERN_ATTR_ESALT (oldoffice34_t))
{
/**
* base
@ -448,7 +448,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_m08 (KERN_A
m09810m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_m16 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09810_m16 (KERN_ATTR_ESALT (oldoffice34_t))
{
/**
* base
@ -497,7 +497,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_m16 (KERN_A
m09810m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_s04 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09810_s04 (KERN_ATTR_ESALT (oldoffice34_t))
{
/**
* base
@ -546,7 +546,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_s04 (KERN_A
m09810s (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_s08 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09810_s08 (KERN_ATTR_ESALT (oldoffice34_t))
{
/**
* base
@ -595,7 +595,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_s08 (KERN_A
m09810s (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_s16 (KERN_ATTR_ESALT (oldoffice34_t))
__kernel void m09810_s16 (KERN_ATTR_ESALT (oldoffice34_t))
{
/**
* base

12
OpenCL/m10400_a0-optimized.cl
View File

@ -157,7 +157,7 @@ DECLSPEC u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __constant u32 *i
return j;
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_m04 (KERN_ATTR_RULES_ESALT (pdf_t))
__kernel void m10400_m04 (KERN_ATTR_RULES_ESALT (pdf_t))
{
/**
* modifier
@ -325,15 +325,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_m08 (KERN_ATTR_RULES_ESALT (pdf_t))
__kernel void m10400_m08 (KERN_ATTR_RULES_ESALT (pdf_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_m16 (KERN_ATTR_RULES_ESALT (pdf_t))
__kernel void m10400_m16 (KERN_ATTR_RULES_ESALT (pdf_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_s04 (KERN_ATTR_RULES_ESALT (pdf_t))
__kernel void m10400_s04 (KERN_ATTR_RULES_ESALT (pdf_t))
{
/**
* modifier
@ -513,10 +513,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_s08 (KERN_ATTR_RULES_ESALT (pdf_t))
__kernel void m10400_s08 (KERN_ATTR_RULES_ESALT (pdf_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_s16 (KERN_ATTR_RULES_ESALT (pdf_t))
__kernel void m10400_s16 (KERN_ATTR_RULES_ESALT (pdf_t))
{
}

12
OpenCL/m10400_a1-optimized.cl
View File

@ -155,7 +155,7 @@ DECLSPEC u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __constant u32 *i
return j;
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_m04 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10400_m04 (KERN_ATTR_ESALT (pdf_t))
{
/**
* modifier
@ -383,15 +383,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_m08 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10400_m08 (KERN_ATTR_ESALT (pdf_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_m16 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10400_m16 (KERN_ATTR_ESALT (pdf_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_s04 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10400_s04 (KERN_ATTR_ESALT (pdf_t))
{
/**
* modifier
@ -631,10 +631,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_s08 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10400_s08 (KERN_ATTR_ESALT (pdf_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_s16 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10400_s16 (KERN_ATTR_ESALT (pdf_t))
{
}

12
OpenCL/m10400_a3-optimized.cl
View File

@ -499,7 +499,7 @@ DECLSPEC void m10400s (__local RC4_KEY *rc4_keys, u32 *w0, u32 *w1, u32 *w2, u32
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_m04 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10400_m04 (KERN_ATTR_ESALT (pdf_t))
{
/**
* base
@ -548,7 +548,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_m04 (KERN_A
m10400m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_m08 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10400_m08 (KERN_ATTR_ESALT (pdf_t))
{
/**
* base
@ -597,7 +597,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_m08 (KERN_A
m10400m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_m16 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10400_m16 (KERN_ATTR_ESALT (pdf_t))
{
/**
* base
@ -646,7 +646,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_m16 (KERN_A
m10400m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_s04 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10400_s04 (KERN_ATTR_ESALT (pdf_t))
{
/**
* base
@ -695,7 +695,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_s04 (KERN_A
m10400s (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_s08 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10400_s08 (KERN_ATTR_ESALT (pdf_t))
{
/**
* base
@ -744,7 +744,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_s08 (KERN_A
m10400s (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_s16 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10400_s16 (KERN_ATTR_ESALT (pdf_t))
{
/**
* base

12
OpenCL/m10410_a0-optimized.cl
View File

@ -157,7 +157,7 @@ DECLSPEC u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __constant u32 *i
return j;
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_m04 (KERN_ATTR_RULES_ESALT (pdf_t))
__kernel void m10410_m04 (KERN_ATTR_RULES_ESALT (pdf_t))
{
/**
* modifier
@ -221,15 +221,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_m08 (KERN_ATTR_RULES_ESALT (pdf_t))
__kernel void m10410_m08 (KERN_ATTR_RULES_ESALT (pdf_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_m16 (KERN_ATTR_RULES_ESALT (pdf_t))
__kernel void m10410_m16 (KERN_ATTR_RULES_ESALT (pdf_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_s04 (KERN_ATTR_RULES_ESALT (pdf_t))
__kernel void m10410_s04 (KERN_ATTR_RULES_ESALT (pdf_t))
{
/**
* modifier
@ -305,10 +305,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_s08 (KERN_ATTR_RULES_ESALT (pdf_t))
__kernel void m10410_s08 (KERN_ATTR_RULES_ESALT (pdf_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_s16 (KERN_ATTR_RULES_ESALT (pdf_t))
__kernel void m10410_s16 (KERN_ATTR_RULES_ESALT (pdf_t))
{
}

12
OpenCL/m10410_a1-optimized.cl
View File

@ -155,7 +155,7 @@ DECLSPEC u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __constant u32 *i
return j;
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_m04 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10410_m04 (KERN_ATTR_ESALT (pdf_t))
{
/**
* modifier
@ -289,15 +289,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_m08 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10410_m08 (KERN_ATTR_ESALT (pdf_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_m16 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10410_m16 (KERN_ATTR_ESALT (pdf_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_s04 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10410_s04 (KERN_ATTR_ESALT (pdf_t))
{
/**
* modifier
@ -443,10 +443,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_s08 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10410_s08 (KERN_ATTR_ESALT (pdf_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_s16 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10410_s16 (KERN_ATTR_ESALT (pdf_t))
{
}

12
OpenCL/m10410_a3-optimized.cl
View File

@ -253,7 +253,7 @@ DECLSPEC void m10410s (__local RC4_KEY *rc4_keys, u32 *w0, u32 *w1, u32 *w2, u32
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_m04 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10410_m04 (KERN_ATTR_ESALT (pdf_t))
{
/**
* base
@ -302,7 +302,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_m04 (KERN_A
m10410m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_m08 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10410_m08 (KERN_ATTR_ESALT (pdf_t))
{
/**
* base
@ -351,7 +351,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_m08 (KERN_A
m10410m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_m16 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10410_m16 (KERN_ATTR_ESALT (pdf_t))
{
/**
* base
@ -400,7 +400,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_m16 (KERN_A
m10410m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_s04 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10410_s04 (KERN_ATTR_ESALT (pdf_t))
{
/**
* base
@ -449,7 +449,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_s04 (KERN_A
m10410s (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_s08 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10410_s08 (KERN_ATTR_ESALT (pdf_t))
{
/**
* base
@ -498,7 +498,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_s08 (KERN_A
m10410s (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_s16 (KERN_ATTR_ESALT (pdf_t))
__kernel void m10410_s16 (KERN_ATTR_ESALT (pdf_t))
{
/**
* base

6
OpenCL/m10500-pure.cl
View File

@ -179,7 +179,7 @@ DECLSPEC u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, const u32 *in, u3
return j;
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10500_init (KERN_ATTR_TMPS_ESALT (pdf14_tmp_t, pdf_t))
__kernel void m10500_init (KERN_ATTR_TMPS_ESALT (pdf14_tmp_t, pdf_t))
{
/**
* base
@ -361,7 +361,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10500_init (KERN_
tmps[gid].out[3] = 0;
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10500_loop (KERN_ATTR_TMPS_ESALT (pdf14_tmp_t, pdf_t))
__kernel void m10500_loop (KERN_ATTR_TMPS_ESALT (pdf14_tmp_t, pdf_t))
{
/**
* base
@ -464,7 +464,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10500_loop (KERN_
tmps[gid].out[3] = out[3];
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10500_comp (KERN_ATTR_TMPS_ESALT (pdf14_tmp_t, pdf_t))
__kernel void m10500_comp (KERN_ATTR_TMPS_ESALT (pdf14_tmp_t, pdf_t))
{
/**
* modifier

12
OpenCL/m13100_a0-optimized.cl
View File

@ -574,7 +574,7 @@ DECLSPEC void kerb_prepare (const u32 *w0, const u32 *w1, const u32 pw_len, cons
hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_m04 (KERN_ATTR_RULES_ESALT (krb5tgs_t))
__kernel void m13100_m04 (KERN_ATTR_RULES_ESALT (krb5tgs_t))
{
/**
* modifier
@ -664,15 +664,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_m08 (KERN_ATTR_RULES_ESALT (krb5tgs_t))
__kernel void m13100_m08 (KERN_ATTR_RULES_ESALT (krb5tgs_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_m16 (KERN_ATTR_RULES_ESALT (krb5tgs_t))
__kernel void m13100_m16 (KERN_ATTR_RULES_ESALT (krb5tgs_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_s04 (KERN_ATTR_RULES_ESALT (krb5tgs_t))
__kernel void m13100_s04 (KERN_ATTR_RULES_ESALT (krb5tgs_t))
{
/**
* modifier
@ -762,10 +762,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_s08 (KERN_ATTR_RULES_ESALT (krb5tgs_t))
__kernel void m13100_s08 (KERN_ATTR_RULES_ESALT (krb5tgs_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_s16 (KERN_ATTR_RULES_ESALT (krb5tgs_t))
__kernel void m13100_s16 (KERN_ATTR_RULES_ESALT (krb5tgs_t))
{
}

4
OpenCL/m13100_a0-pure.cl
View File

@ -386,7 +386,7 @@ DECLSPEC void kerb_prepare (const u32 *K, const u32 *checksum, u32 *digest, u32
K2[3] = ctx1.opad.h[3];
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_mxx (KERN_ATTR_RULES_ESALT (krb5tgs_t))
__kernel void m13100_mxx (KERN_ATTR_RULES_ESALT (krb5tgs_t))
{
/**
* modifier
@ -448,7 +448,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_mxx (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_sxx (KERN_ATTR_RULES_ESALT (krb5tgs_t))
__kernel void m13100_sxx (KERN_ATTR_RULES_ESALT (krb5tgs_t))
{
/**
* modifier

12
OpenCL/m13100_a1-optimized.cl
View File

@ -572,7 +572,7 @@ DECLSPEC void kerb_prepare (const u32 *w0, const u32 *w1, const u32 pw_len, cons
hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_m04 (KERN_ATTR_ESALT (krb5tgs_t))
__kernel void m13100_m04 (KERN_ATTR_ESALT (krb5tgs_t))
{
/**
* modifier
@ -711,15 +711,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_m08 (KERN_ATTR_ESALT (krb5tgs_t))
__kernel void m13100_m08 (KERN_ATTR_ESALT (krb5tgs_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_m16 (KERN_ATTR_ESALT (krb5tgs_t))
__kernel void m13100_m16 (KERN_ATTR_ESALT (krb5tgs_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_s04 (KERN_ATTR_ESALT (krb5tgs_t))
__kernel void m13100_s04 (KERN_ATTR_ESALT (krb5tgs_t))
{
/**
* modifier
@ -858,10 +858,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_s08 (KERN_ATTR_ESALT (krb5tgs_t))
__kernel void m13100_s08 (KERN_ATTR_ESALT (krb5tgs_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_s16 (KERN_ATTR_ESALT (krb5tgs_t))
__kernel void m13100_s16 (KERN_ATTR_ESALT (krb5tgs_t))
{
}

4
OpenCL/m13100_a1-pure.cl
View File

@ -384,7 +384,7 @@ DECLSPEC void kerb_prepare (const u32 *K, const u32 *checksum, u32 *digest, u32
K2[3] = ctx1.opad.h[3];
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_mxx (KERN_ATTR_ESALT (krb5tgs_t))
__kernel void m13100_mxx (KERN_ATTR_ESALT (krb5tgs_t))
{
/**
* modifier
@ -444,7 +444,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_mxx (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_sxx (KERN_ATTR_ESALT (krb5tgs_t))
__kernel void m13100_sxx (KERN_ATTR_ESALT (krb5tgs_t))
{
/**
* modifier

12
OpenCL/m13100_a3-optimized.cl
View File

@ -631,7 +631,7 @@ DECLSPEC void m13100 (__local RC4_KEY *rc4_key, u32 *w0, u32 *w1, u32 *w2, u32 *
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_m04 (KERN_ATTR_ESALT (krb5tgs_t))
__kernel void m13100_m04 (KERN_ATTR_ESALT (krb5tgs_t))
{
/**
* base
@ -683,7 +683,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_m04 (KERN_A
m13100 (rc4_key, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_m08 (KERN_ATTR_ESALT (krb5tgs_t))
__kernel void m13100_m08 (KERN_ATTR_ESALT (krb5tgs_t))
{
/**
* base
@ -735,11 +735,11 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_m08 (KERN_A
m13100 (rc4_key, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_m16 (KERN_ATTR_ESALT (krb5tgs_t))
__kernel void m13100_m16 (KERN_ATTR_ESALT (krb5tgs_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_s04 (KERN_ATTR_ESALT (krb5tgs_t))
__kernel void m13100_s04 (KERN_ATTR_ESALT (krb5tgs_t))
{
/**
* base
@ -791,7 +791,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_s04 (KERN_A
m13100 (rc4_key, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_s08 (KERN_ATTR_ESALT (krb5tgs_t))
__kernel void m13100_s08 (KERN_ATTR_ESALT (krb5tgs_t))
{
/**
* base
@ -843,6 +843,6 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_s08 (KERN_A
m13100 (rc4_key, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_s16 (KERN_ATTR_ESALT (krb5tgs_t))
__kernel void m13100_s16 (KERN_ATTR_ESALT (krb5tgs_t))
{
}

4
OpenCL/m13100_a3-pure.cl
View File

@ -384,7 +384,7 @@ DECLSPEC void kerb_prepare (const u32 *K, const u32 *checksum, u32 *digest, u32
K2[3] = ctx1.opad.h[3];
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_mxx (KERN_ATTR_VECTOR_ESALT (krb5tgs_t))
__kernel void m13100_mxx (KERN_ATTR_VECTOR_ESALT (krb5tgs_t))
{
/**
* modifier
@ -457,7 +457,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_mxx (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_sxx (KERN_ATTR_VECTOR_ESALT (krb5tgs_t))
__kernel void m13100_sxx (KERN_ATTR_VECTOR_ESALT (krb5tgs_t))
{
/**
* modifier

6
OpenCL/m15700-pure.cl
View File

@ -353,7 +353,7 @@ DECLSPEC void keccak_transform_S (u64 *st)
}
}
__kernel void __attribute__((reqd_work_group_size(1, 1, 1))) m15700_init (KERN_ATTR_TMPS_ESALT (scrypt_tmp_t, ethereum_scrypt_t))
__kernel void m15700_init (KERN_ATTR_TMPS_ESALT (scrypt_tmp_t, ethereum_scrypt_t))
{
/**
* base
@ -418,7 +418,7 @@ __kernel void __attribute__((reqd_work_group_size(1, 1, 1))) m15700_init (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(1, 1, 1))) m15700_loop (KERN_ATTR_TMPS_ESALT (scrypt_tmp_t, ethereum_scrypt_t))
__kernel void m15700_loop (KERN_ATTR_TMPS_ESALT (scrypt_tmp_t, ethereum_scrypt_t))
{
const u64 gid = get_global_id (0);
@ -456,7 +456,7 @@ __kernel void __attribute__((reqd_work_group_size(1, 1, 1))) m15700_loop (KERN_A
#endif
}
__kernel void __attribute__((reqd_work_group_size(1, 1, 1))) m15700_comp (KERN_ATTR_TMPS_ESALT (scrypt_tmp_t, ethereum_scrypt_t))
__kernel void m15700_comp (KERN_ATTR_TMPS_ESALT (scrypt_tmp_t, ethereum_scrypt_t))
{
/**
* base

12
OpenCL/m18200_a0-optimized.cl
View File

@ -572,7 +572,7 @@ DECLSPEC void kerb_prepare (const u32 *w0, const u32 *w1, const u32 pw_len, cons
hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_m04 (KERN_ATTR_RULES_ESALT (krb5asrep_t))
__kernel void m18200_m04 (KERN_ATTR_RULES_ESALT (krb5asrep_t))
{
/**
* modifier
@ -662,15 +662,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_m08 (KERN_ATTR_RULES_ESALT (krb5asrep_t))
__kernel void m18200_m08 (KERN_ATTR_RULES_ESALT (krb5asrep_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_m16 (KERN_ATTR_RULES_ESALT (krb5asrep_t))
__kernel void m18200_m16 (KERN_ATTR_RULES_ESALT (krb5asrep_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_s04 (KERN_ATTR_RULES_ESALT (krb5asrep_t))
__kernel void m18200_s04 (KERN_ATTR_RULES_ESALT (krb5asrep_t))
{
/**
* modifier
@ -760,10 +760,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_s08 (KERN_ATTR_RULES_ESALT (krb5asrep_t))
__kernel void m18200_s08 (KERN_ATTR_RULES_ESALT (krb5asrep_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_s16 (KERN_ATTR_RULES_ESALT (krb5asrep_t))
__kernel void m18200_s16 (KERN_ATTR_RULES_ESALT (krb5asrep_t))
{
}

4
OpenCL/m18200_a0-pure.cl
View File

@ -384,7 +384,7 @@ DECLSPEC void kerb_prepare (const u32 *K, const u32 *checksum, u32 *digest, u32
K2[3] = ctx1.opad.h[3];
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_mxx (KERN_ATTR_RULES_ESALT (krb5asrep_t))
__kernel void m18200_mxx (KERN_ATTR_RULES_ESALT (krb5asrep_t))
{
/**
* modifier
@ -446,7 +446,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_mxx (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_sxx (KERN_ATTR_RULES_ESALT (krb5asrep_t))
__kernel void m18200_sxx (KERN_ATTR_RULES_ESALT (krb5asrep_t))
{
/**
* modifier

12
OpenCL/m18200_a1-optimized.cl
View File

@ -570,7 +570,7 @@ DECLSPEC void kerb_prepare (const u32 *w0, const u32 *w1, const u32 pw_len, cons
hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_m04 (KERN_ATTR_ESALT (krb5asrep_t))
__kernel void m18200_m04 (KERN_ATTR_ESALT (krb5asrep_t))
{
/**
* modifier
@ -709,15 +709,15 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_m04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_m08 (KERN_ATTR_ESALT (krb5asrep_t))
__kernel void m18200_m08 (KERN_ATTR_ESALT (krb5asrep_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_m16 (KERN_ATTR_ESALT (krb5asrep_t))
__kernel void m18200_m16 (KERN_ATTR_ESALT (krb5asrep_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_s04 (KERN_ATTR_ESALT (krb5asrep_t))
__kernel void m18200_s04 (KERN_ATTR_ESALT (krb5asrep_t))
{
/**
* modifier
@ -856,10 +856,10 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_s04 (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_s08 (KERN_ATTR_ESALT (krb5asrep_t))
__kernel void m18200_s08 (KERN_ATTR_ESALT (krb5asrep_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_s16 (KERN_ATTR_ESALT (krb5asrep_t))
__kernel void m18200_s16 (KERN_ATTR_ESALT (krb5asrep_t))
{
}

4
OpenCL/m18200_a1-pure.cl
View File

@ -382,7 +382,7 @@ DECLSPEC void kerb_prepare (const u32 *K, const u32 *checksum, u32 *digest, u32
K2[3] = ctx1.opad.h[3];
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_mxx (KERN_ATTR_ESALT (krb5asrep_t))
__kernel void m18200_mxx (KERN_ATTR_ESALT (krb5asrep_t))
{
/**
* modifier
@ -442,7 +442,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_mxx (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_sxx (KERN_ATTR_ESALT (krb5asrep_t))
__kernel void m18200_sxx (KERN_ATTR_ESALT (krb5asrep_t))
{
/**
* modifier

12
OpenCL/m18200_a3-optimized.cl
View File

@ -629,7 +629,7 @@ DECLSPEC void m18200 (__local RC4_KEY *rc4_key, u32 *w0, u32 *w1, u32 *w2, u32 *
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_m04 (KERN_ATTR_ESALT (krb5asrep_t))
__kernel void m18200_m04 (KERN_ATTR_ESALT (krb5asrep_t))
{
/**
* base
@ -681,7 +681,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_m04 (KERN_A
m18200 (rc4_key, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_m08 (KERN_ATTR_ESALT (krb5asrep_t))
__kernel void m18200_m08 (KERN_ATTR_ESALT (krb5asrep_t))
{
/**
* base
@ -733,11 +733,11 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_m08 (KERN_A
m18200 (rc4_key, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_m16 (KERN_ATTR_ESALT (krb5asrep_t))
__kernel void m18200_m16 (KERN_ATTR_ESALT (krb5asrep_t))
{
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_s04 (KERN_ATTR_ESALT (krb5asrep_t))
__kernel void m18200_s04 (KERN_ATTR_ESALT (krb5asrep_t))
{
/**
* base
@ -789,7 +789,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_s04 (KERN_A
m18200 (rc4_key, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_s08 (KERN_ATTR_ESALT (krb5asrep_t))
__kernel void m18200_s08 (KERN_ATTR_ESALT (krb5asrep_t))
{
/**
* base
@ -841,6 +841,6 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_s08 (KERN_A
m18200 (rc4_key, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_s16 (KERN_ATTR_ESALT (krb5asrep_t))
__kernel void m18200_s16 (KERN_ATTR_ESALT (krb5asrep_t))
{
}

4
OpenCL/m18200_a3-pure.cl
View File

@ -382,7 +382,7 @@ DECLSPEC void kerb_prepare (const u32 *K, const u32 *checksum, u32 *digest, u32
K2[3] = ctx1.opad.h[3];
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_mxx (KERN_ATTR_VECTOR_ESALT (krb5asrep_t))
__kernel void m18200_mxx (KERN_ATTR_VECTOR_ESALT (krb5asrep_t))
{
/**
* modifier
@ -455,7 +455,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_mxx (KERN_A
}
}
__kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_sxx (KERN_ATTR_VECTOR_ESALT (krb5asrep_t))
__kernel void m18200_sxx (KERN_ATTR_VECTOR_ESALT (krb5asrep_t))
{
/**
* modifier

2
OpenCL/m18600-pure.cl
View File

@ -582,7 +582,7 @@ __kernel void m18600_loop (KERN_ATTR_TMPS_ESALT (odf11_tmp_t, odf11_t))
}
}
__kernel void __attribute__((reqd_work_group_size(8, 1, 1))) m18600_comp (KERN_ATTR_TMPS_ESALT (odf11_tmp_t, odf11_t))
__kernel void m18600_comp (KERN_ATTR_TMPS_ESALT (odf11_tmp_t, odf11_t))
{
const u64 gid = get_global_id (0);
const u64 lid = get_local_id (0);

16
src/modules/module_03200.c
View File

@ -86,20 +86,6 @@ u32 module_kernel_threads_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYB
return kernel_threads_max;
}
bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hc_device_param_t *device_param)
{
// CL_INVALID_WORK_GROUP_SIZE
if (device_param->platform_vendor_id == VENDOR_ID_APPLE)
{
if (device_param->device_type & CL_DEVICE_TYPE_CPU)
{
return true;
}
}
return false;
}
int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len)
{
u32 *digest = (u32 *) digest_buf;
@ -276,6 +262,6 @@ void module_init (module_ctx_t *module_ctx)
module_ctx->module_st_hash = module_st_hash;
module_ctx->module_st_pass = module_st_pass;
module_ctx->module_tmp_size = module_tmp_size;
module_ctx->module_unstable_warning = module_unstable_warning;
module_ctx->module_unstable_warning = MODULE_DEFAULT;
module_ctx->module_warmup_disable = MODULE_DEFAULT;
}

28
src/modules/module_07500.c
View File

@ -76,32 +76,6 @@ u32 module_kernel_threads_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYB
return kernel_threads_max;
}
bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hc_device_param_t *device_param)
{
// amdgpu-pro-18.50-708488-ubuntu-18.04: CL_OUT_OF_RESOURCES
if ((device_param->device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == false))
{
if ((hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) == 0)
{
return true;
}
}
// CL_INVALID_WORK_GROUP_SIZE
if (device_param->platform_vendor_id == VENDOR_ID_APPLE)
{
if (device_param->device_type & CL_DEVICE_TYPE_CPU)
{
if ((hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) == 1)
{
return true;
}
}
}
return false;
}
int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len)
{
u32 *digest = (u32 *) digest_buf;
@ -316,6 +290,6 @@ void module_init (module_ctx_t *module_ctx)
module_ctx->module_st_hash = module_st_hash;
module_ctx->module_st_pass = module_st_pass;
module_ctx->module_tmp_size = MODULE_DEFAULT;
module_ctx->module_unstable_warning = module_unstable_warning;
module_ctx->module_unstable_warning = MODULE_DEFAULT;
module_ctx->module_warmup_disable = MODULE_DEFAULT;
}

28
src/modules/module_13100.c
View File

@ -75,32 +75,6 @@ u32 module_kernel_threads_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYB
return kernel_threads_max;
}
bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hc_device_param_t *device_param)
{
// amdgpu-pro-18.50-708488-ubuntu-18.04: CL_OUT_OF_RESOURCES
if ((device_param->device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == false))
{
if ((hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) == 0)
{
return true;
}
}
// CL_INVALID_WORK_GROUP_SIZE
if (device_param->platform_vendor_id == VENDOR_ID_APPLE)
{
if (device_param->device_type & CL_DEVICE_TYPE_CPU)
{
if ((hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) == 1)
{
return true;
}
}
}
return false;
}
int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len)
{
u32 *digest = (u32 *) digest_buf;
@ -325,6 +299,6 @@ void module_init (module_ctx_t *module_ctx)
module_ctx->module_st_hash = module_st_hash;
module_ctx->module_st_pass = module_st_pass;
module_ctx->module_tmp_size = MODULE_DEFAULT;
module_ctx->module_unstable_warning = module_unstable_warning;
module_ctx->module_unstable_warning = MODULE_DEFAULT;
module_ctx->module_warmup_disable = MODULE_DEFAULT;
}

16
src/modules/module_18200.c
View File

@ -75,20 +75,6 @@ u32 module_kernel_threads_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYB
return kernel_threads_max;
}
bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hc_device_param_t *hc_device_param)
{
// amdgpu-pro-18.50-708488-ubuntu-18.04: CL_OUT_OF_RESOURCES
if ((hc_device_param->device_vendor_id == VENDOR_ID_AMD) && (hc_device_param->has_vperm == false))
{
if ((hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) == 0)
{
return true;
}
}
return false;
}
int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len)
{
u32 *digest = (u32 *) digest_buf;
@ -275,6 +261,6 @@ void module_init (module_ctx_t *module_ctx)
module_ctx->module_st_hash = module_st_hash;
module_ctx->module_st_pass = module_st_pass;
module_ctx->module_tmp_size = MODULE_DEFAULT;
module_ctx->module_unstable_warning = module_unstable_warning;
module_ctx->module_unstable_warning = MODULE_DEFAULT;
module_ctx->module_warmup_disable = MODULE_DEFAULT;
}