arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-22 16:18:09 +00:00
Code Issues Releases Wiki Activity

Do not use __local memory for -m 13100 if running on a device without physical shared memory

Browse Source
This commit is contained in:
Jens Steube 2017-09-05 17:41:55 +02:00
parent b58aa445b4
commit a910aea9e0
6 changed files with 214 additions and 46 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

38
OpenCL/m13100_a0-optimized.cl
View File

@ -25,7 +25,7 @@ typedef struct
} RC4_KEY;
void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
void swap (SCR_TYPE RC4_KEY *rc4_key, const u8 i, const u8 j)
{
u8 tmp;
@ -34,12 +34,12 @@ void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
rc4_key->S[j] = tmp;
}
void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
void rc4_init_16 (SCR_TYPE RC4_KEY *rc4_key, const u32 data[4])
{
u32 v = 0x03020100;
u32 a = 0x04040404;
__local u32 *ptr = (__local u32 *) rc4_key->S;
SCR_TYPE u32 *ptr = (SCR_TYPE u32 *) rc4_key->S;
#ifdef _unroll
#pragma unroll
@ -87,7 +87,7 @@ void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
}
}
u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __global u32 *in, u32 out[4])
u8 rc4_next_16 (SCR_TYPE RC4_KEY *rc4_key, u8 i, u8 j, __global u32 *in, u32 out[4])
{
#ifdef _unroll
#pragma unroll
@ -225,7 +225,7 @@ void hmac_md5_run (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[4], u32
md5_transform (w0, w1, w2, w3, digest);
}
int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edata2, const u32 edata2_len, const u32 K2[4], const u32 checksum[4])
int decrypt_and_check (SCR_TYPE RC4_KEY *rc4_key, u32 data[4], __global u32 *edata2, const u32 edata2_len, const u32 K2[4], const u32 checksum[4])
{
rc4_init_16 (rc4_key, data);
@ -600,8 +600,20 @@ __kernel void m13100_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru
* shared
*/
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64];
__local RC4_KEY *rc4_key = &rc4_keys[lid];
#else
RC4_KEY rc4_keys[1];
RC4_KEY *rc4_key = &rc4_keys[0];
#endif
/**
* salt
*/
@ -643,7 +655,7 @@ __kernel void m13100_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru
tmp[2] = digest[2];
tmp[3] = digest[3];
if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
if (decrypt_and_check (rc4_key, tmp, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
@ -696,8 +708,20 @@ __kernel void m13100_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru
* shared
*/
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64];
__local RC4_KEY *rc4_key = &rc4_keys[lid];
#else
RC4_KEY rc4_keys[1];
RC4_KEY *rc4_key = &rc4_keys[0];
#endif
/**
* salt
*/
@ -739,7 +763,7 @@ __kernel void m13100_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru
tmp[2] = digest[2];
tmp[3] = digest[3];
if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
if (decrypt_and_check (rc4_key, tmp, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{

38
OpenCL/m13100_a0.cl
View File

@ -24,7 +24,7 @@ typedef struct
} RC4_KEY;
void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
void swap (SCR_TYPE RC4_KEY *rc4_key, const u8 i, const u8 j)
{
u8 tmp;
@ -33,12 +33,12 @@ void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
rc4_key->S[j] = tmp;
}
void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
void rc4_init_16 (SCR_TYPE RC4_KEY *rc4_key, const u32 data[4])
{
u32 v = 0x03020100;
u32 a = 0x04040404;
__local u32 *ptr = (__local u32 *) rc4_key->S;
SCR_TYPE u32 *ptr = (SCR_TYPE u32 *) rc4_key->S;
#ifdef _unroll
#pragma unroll
@ -86,7 +86,7 @@ void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
}
}
u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __global const u32 in[4], u32 out[4])
u8 rc4_next_16 (SCR_TYPE RC4_KEY *rc4_key, u8 i, u8 j, __global const u32 in[4], u32 out[4])
{
#ifdef _unroll
#pragma unroll
@ -139,7 +139,7 @@ u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __global const u32 in[4],
return j;
}
int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32 *edata2, const u32 edata2_len, const u32 K2[4], const u32 checksum[4])
int decrypt_and_check (SCR_TYPE RC4_KEY *rc4_key, u32 data[4], __global const u32 *edata2, const u32 edata2_len, const u32 K2[4], const u32 checksum[4])
{
rc4_init_16 (rc4_key, data);
@ -394,8 +394,20 @@ __kernel void m13100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru
COPY_PW (pws[gid]);
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64];
__local RC4_KEY *rc4_key = &rc4_keys[lid];
#else
RC4_KEY rc4_keys[1];
RC4_KEY *rc4_key = &rc4_keys[0];
#endif
u32 checksum[4];
checksum[0] = krb5tgs_bufs[digests_offset].checksum[0];
@ -427,7 +439,7 @@ __kernel void m13100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru
kerb_prepare (ctx.h, checksum, digest, K2);
if (decrypt_and_check (&rc4_keys[lid], digest, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
if (decrypt_and_check (rc4_key, digest, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
@ -454,8 +466,20 @@ __kernel void m13100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru
COPY_PW (pws[gid]);
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64];
__local RC4_KEY *rc4_key = &rc4_keys[lid];
#else
RC4_KEY rc4_keys[1];
RC4_KEY *rc4_key = &rc4_keys[0];
#endif
u32 checksum[4];
checksum[0] = krb5tgs_bufs[digests_offset].checksum[0];
@ -487,7 +511,7 @@ __kernel void m13100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru
kerb_prepare (ctx.h, checksum, digest, K2);
if (decrypt_and_check (&rc4_keys[lid], digest, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
if (decrypt_and_check (rc4_key, digest, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{

38
OpenCL/m13100_a1-optimized.cl
View File

@ -23,7 +23,7 @@ typedef struct
} RC4_KEY;
void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
void swap (SCR_TYPE RC4_KEY *rc4_key, const u8 i, const u8 j)
{
u8 tmp;
@ -32,12 +32,12 @@ void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
rc4_key->S[j] = tmp;
}
void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
void rc4_init_16 (SCR_TYPE RC4_KEY *rc4_key, const u32 data[4])
{
u32 v = 0x03020100;
u32 a = 0x04040404;
__local u32 *ptr = (__local u32 *) rc4_key->S;
SCR_TYPE u32 *ptr = (SCR_TYPE u32 *) rc4_key->S;
#ifdef _unroll
#pragma unroll
@ -85,7 +85,7 @@ void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
}
}
u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __global u32 *in, u32 out[4])
u8 rc4_next_16 (SCR_TYPE RC4_KEY *rc4_key, u8 i, u8 j, __global u32 *in, u32 out[4])
{
#ifdef _unroll
#pragma unroll
@ -223,7 +223,7 @@ void hmac_md5_run (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[4], u32
md5_transform (w0, w1, w2, w3, digest);
}
int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edata2, const u32 edata2_len, const u32 K2[4], const u32 checksum[4])
int decrypt_and_check (SCR_TYPE RC4_KEY *rc4_key, u32 data[4], __global u32 *edata2, const u32 edata2_len, const u32 K2[4], const u32 checksum[4])
{
rc4_init_16 (rc4_key, data);
@ -597,8 +597,20 @@ __kernel void m13100_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* shared
*/
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64];
__local RC4_KEY *rc4_key = &rc4_keys[lid];
#else
RC4_KEY rc4_keys[1];
RC4_KEY *rc4_key = &rc4_keys[0];
#endif
/**
* salt
*/
@ -690,7 +702,7 @@ __kernel void m13100_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
tmp[2] = digest[2];
tmp[3] = digest[3];
if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
if (decrypt_and_check (rc4_key, tmp, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
@ -742,8 +754,20 @@ __kernel void m13100_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* shared
*/
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64];
__local RC4_KEY *rc4_key = &rc4_keys[lid];
#else
RC4_KEY rc4_keys[1];
RC4_KEY *rc4_key = &rc4_keys[0];
#endif
/**
* salt
*/
@ -835,7 +859,7 @@ __kernel void m13100_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
tmp[2] = digest[2];
tmp[3] = digest[3];
if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
if (decrypt_and_check (rc4_key, tmp, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{

38
OpenCL/m13100_a1.cl
View File

@ -22,7 +22,7 @@ typedef struct
} RC4_KEY;
void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
void swap (SCR_TYPE RC4_KEY *rc4_key, const u8 i, const u8 j)
{
u8 tmp;
@ -31,12 +31,12 @@ void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
rc4_key->S[j] = tmp;
}
void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
void rc4_init_16 (SCR_TYPE RC4_KEY *rc4_key, const u32 data[4])
{
u32 v = 0x03020100;
u32 a = 0x04040404;
__local u32 *ptr = (__local u32 *) rc4_key->S;
SCR_TYPE u32 *ptr = (SCR_TYPE u32 *) rc4_key->S;
#ifdef _unroll
#pragma unroll
@ -84,7 +84,7 @@ void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
}
}
u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __global const u32 in[4], u32 out[4])
u8 rc4_next_16 (SCR_TYPE RC4_KEY *rc4_key, u8 i, u8 j, __global const u32 in[4], u32 out[4])
{
#ifdef _unroll
#pragma unroll
@ -137,7 +137,7 @@ u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __global const u32 in[4],
return j;
}
int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32 *edata2, const u32 edata2_len, const u32 K2[4], const u32 checksum[4])
int decrypt_and_check (SCR_TYPE RC4_KEY *rc4_key, u32 data[4], __global const u32 *edata2, const u32 edata2_len, const u32 K2[4], const u32 checksum[4])
{
rc4_init_16 (rc4_key, data);
@ -390,8 +390,20 @@ __kernel void m13100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
* base
*/
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64];
__local RC4_KEY *rc4_key = &rc4_keys[lid];
#else
RC4_KEY rc4_keys[1];
RC4_KEY *rc4_key = &rc4_keys[0];
#endif
u32 checksum[4];
checksum[0] = krb5tgs_bufs[digests_offset].checksum[0];
@ -423,7 +435,7 @@ __kernel void m13100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
kerb_prepare (ctx.h, checksum, digest, K2);
if (decrypt_and_check (&rc4_keys[lid], digest, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
if (decrypt_and_check (rc4_key, digest, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
@ -448,8 +460,20 @@ __kernel void m13100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
* base
*/
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64];
__local RC4_KEY *rc4_key = &rc4_keys[lid];
#else
RC4_KEY rc4_keys[1];
RC4_KEY *rc4_key = &rc4_keys[0];
#endif
u32 checksum[4];
checksum[0] = krb5tgs_bufs[digests_offset].checksum[0];
@ -481,7 +505,7 @@ __kernel void m13100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
kerb_prepare (ctx.h, checksum, digest, K2);
if (decrypt_and_check (&rc4_keys[lid], digest, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
if (decrypt_and_check (rc4_key, digest, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{

70
OpenCL/m13100_a3-optimized.cl
View File

@ -23,7 +23,7 @@ typedef struct
} RC4_KEY;
void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
void swap (SCR_TYPE RC4_KEY *rc4_key, const u8 i, const u8 j)
{
u8 tmp;
@ -32,12 +32,12 @@ void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
rc4_key->S[j] = tmp;
}
void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
void rc4_init_16 (SCR_TYPE RC4_KEY *rc4_key, const u32 data[4])
{
u32 v = 0x03020100;
u32 a = 0x04040404;
__local u32 *ptr = (__local u32 *) rc4_key->S;
SCR_TYPE u32 *ptr = (SCR_TYPE u32 *) rc4_key->S;
#ifdef _unroll
#pragma unroll
@ -85,7 +85,7 @@ void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
}
}
u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __global u32 *in, u32 out[4])
u8 rc4_next_16 (SCR_TYPE RC4_KEY *rc4_key, u8 i, u8 j, __global u32 *in, u32 out[4])
{
#ifdef _unroll
#pragma unroll
@ -223,7 +223,7 @@ void hmac_md5_run (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[4], u32
md5_transform (w0, w1, w2, w3, digest);
}
int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edata2, const u32 edata2_len, const u32 K2[4], const u32 checksum[4])
int decrypt_and_check (SCR_TYPE RC4_KEY *rc4_key, u32 data[4], __global u32 *edata2, const u32 edata2_len, const u32 K2[4], const u32 checksum[4])
{
rc4_init_16 (rc4_key, data);
@ -563,7 +563,7 @@ void kerb_prepare (const u32 w0[4], const u32 w1[4], const u32 pw_len, const u32
hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
}
void m13100 (__local RC4_KEY *rc4_keys, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __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 krb5tgs_t *krb5tgs_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 m13100 (SCR_TYPE RC4_KEY *rc4_key, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __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 krb5tgs_t *krb5tgs_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
@ -612,7 +612,7 @@ void m13100 (__local RC4_KEY *rc4_keys, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[
tmp[2] = digest[2];
tmp[3] = digest[3];
if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
if (decrypt_and_check (rc4_key, tmp, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
@ -667,9 +667,21 @@ __kernel void m13100_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* main
*/
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64];
m13100 (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, krb5tgs_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);
__local RC4_KEY *rc4_key = &rc4_keys[lid];
#else
RC4_KEY rc4_keys[1];
RC4_KEY *rc4_key = &rc4_keys[0];
#endif
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, krb5tgs_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 m13100_m08 (__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 krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
@ -717,9 +729,21 @@ __kernel void m13100_m08 (__global pw_t *pws, __global const kernel_rule_t *rule
* main
*/
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64];
m13100 (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, krb5tgs_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);
__local RC4_KEY *rc4_key = &rc4_keys[lid];
#else
RC4_KEY rc4_keys[1];
RC4_KEY *rc4_key = &rc4_keys[0];
#endif
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, krb5tgs_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 m13100_m16 (__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 krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
@ -771,9 +795,21 @@ __kernel void m13100_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* main
*/
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64];
m13100 (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, krb5tgs_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);
__local RC4_KEY *rc4_key = &rc4_keys[lid];
#else
RC4_KEY rc4_keys[1];
RC4_KEY *rc4_key = &rc4_keys[0];
#endif
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, krb5tgs_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 m13100_s08 (__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 krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
@ -821,9 +857,21 @@ __kernel void m13100_s08 (__global pw_t *pws, __global const kernel_rule_t *rule
* main
*/
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64];
m13100 (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, krb5tgs_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);
__local RC4_KEY *rc4_key = &rc4_keys[lid];
#else
RC4_KEY rc4_keys[1];
RC4_KEY *rc4_key = &rc4_keys[0];
#endif
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, krb5tgs_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 m13100_s16 (__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 krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)

38
OpenCL/m13100_a3.cl
View File

@ -22,7 +22,7 @@ typedef struct
} RC4_KEY;
void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
void swap (SCR_TYPE RC4_KEY *rc4_key, const u8 i, const u8 j)
{
u8 tmp;
@ -31,12 +31,12 @@ void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
rc4_key->S[j] = tmp;
}
void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
void rc4_init_16 (SCR_TYPE RC4_KEY *rc4_key, const u32 data[4])
{
u32 v = 0x03020100;
u32 a = 0x04040404;
__local u32 *ptr = (__local u32 *) rc4_key->S;
SCR_TYPE u32 *ptr = (SCR_TYPE u32 *) rc4_key->S;
#ifdef _unroll
#pragma unroll
@ -84,7 +84,7 @@ void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
}
}
u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __global const u32 in[4], u32 out[4])
u8 rc4_next_16 (SCR_TYPE RC4_KEY *rc4_key, u8 i, u8 j, __global const u32 in[4], u32 out[4])
{
#ifdef _unroll
#pragma unroll
@ -137,7 +137,7 @@ u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __global const u32 in[4],
return j;
}
int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32 *edata2, const u32 edata2_len, const u32 K2[4], const u32 checksum[4])
int decrypt_and_check (SCR_TYPE RC4_KEY *rc4_key, u32 data[4], __global const u32 *edata2, const u32 edata2_len, const u32 K2[4], const u32 checksum[4])
{
rc4_init_16 (rc4_key, data);
@ -399,8 +399,20 @@ __kernel void m13100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
w[idx] = pws[gid].i[idx];
}
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64];
__local RC4_KEY *rc4_key = &rc4_keys[lid];
#else
RC4_KEY rc4_keys[1];
RC4_KEY *rc4_key = &rc4_keys[0];
#endif
u32 checksum[4];
checksum[0] = krb5tgs_bufs[digests_offset].checksum[0];
@ -436,7 +448,7 @@ __kernel void m13100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
kerb_prepare (ctx.h, checksum, digest, K2);
if (decrypt_and_check (&rc4_keys[lid], digest, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
if (decrypt_and_check (rc4_key, digest, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
@ -470,8 +482,20 @@ __kernel void m13100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
w[idx] = pws[gid].i[idx];
}
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64];
__local RC4_KEY *rc4_key = &rc4_keys[lid];
#else
RC4_KEY rc4_keys[1];
RC4_KEY *rc4_key = &rc4_keys[0];
#endif
u32 checksum[4];
checksum[0] = krb5tgs_bufs[digests_offset].checksum[0];
@ -507,7 +531,7 @@ __kernel void m13100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
kerb_prepare (ctx.h, checksum, digest, K2);
if (decrypt_and_check (&rc4_keys[lid], digest, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
if (decrypt_and_check (rc4_key, digest, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{