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

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

Browse Source
This commit is contained in:
Jens Steube 2017-09-05 17:37:20 +02:00
parent 3e597f4c7b
commit b58aa445b4
7 changed files with 216 additions and 46 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
OpenCL/inc_vendor.cl
View File

@ -27,8 +27,10 @@
#ifdef REAL_SHM #ifdef REAL_SHM
#define SHM_TYPE __local #define SHM_TYPE __local
#define SCR_TYPE __local
#else #else
#define SHM_TYPE __constant #define SHM_TYPE __constant
#define SCR_TYPE
#endif #endif
/** /**

38
OpenCL/m07500_a0-optimized.cl
View File

@ -25,7 +25,7 @@ typedef struct
} RC4_KEY; } 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; u8 tmp;
@ -34,12 +34,12 @@ void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
rc4_key->S[j] = tmp; 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 v = 0x03020100;
u32 a = 0x04040404; u32 a = 0x04040404;
__local u32 *ptr = (__local u32 *) rc4_key->S; SCR_TYPE u32 *ptr = (SCR_TYPE u32 *) rc4_key->S;
#ifdef _unroll #ifdef _unroll
#pragma 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, const u32 in[4], u32 out[4]) u8 rc4_next_16 (SCR_TYPE RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4])
{ {
#ifdef _unroll #ifdef _unroll
#pragma unroll #pragma unroll
@ -140,7 +140,7 @@ u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4
return j; return j;
} }
int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], u32 timestamp_ct[8]) int decrypt_and_check (SCR_TYPE RC4_KEY *rc4_key, u32 data[4], u32 timestamp_ct[8])
{ {
rc4_init_16 (rc4_key, data); rc4_init_16 (rc4_key, data);
@ -440,8 +440,20 @@ __kernel void m07500_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru
* shared * shared
*/ */
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64]; __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
/** /**
* loop * loop
*/ */
@ -470,7 +482,7 @@ __kernel void m07500_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru
tmp[2] = digest[2]; tmp[2] = digest[2];
tmp[3] = digest[3]; tmp[3] = digest[3];
if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1) if (decrypt_and_check (rc4_key, tmp, timestamp_ct) == 1)
{ {
if (atomic_inc (&hashes_shown[digests_offset]) == 0) if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{ {
@ -544,8 +556,20 @@ __kernel void m07500_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru
* shared * shared
*/ */
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64]; __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
/** /**
* loop * loop
*/ */
@ -574,7 +598,7 @@ __kernel void m07500_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru
tmp[2] = digest[2]; tmp[2] = digest[2];
tmp[3] = digest[3]; tmp[3] = digest[3];
if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1) if (decrypt_and_check (rc4_key, tmp, timestamp_ct) == 1)
{ {
if (atomic_inc (&hashes_shown[digests_offset]) == 0) if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{ {

38
OpenCL/m07500_a0.cl
View File

@ -24,7 +24,7 @@ typedef struct
} RC4_KEY; } 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; u8 tmp;
@ -33,12 +33,12 @@ void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
rc4_key->S[j] = tmp; 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 v = 0x03020100;
u32 a = 0x04040404; u32 a = 0x04040404;
__local u32 *ptr = (__local u32 *) rc4_key->S; SCR_TYPE u32 *ptr = (SCR_TYPE u32 *) rc4_key->S;
#ifdef _unroll #ifdef _unroll
#pragma 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, const u32 in[4], u32 out[4]) u8 rc4_next_16 (SCR_TYPE RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4])
{ {
#ifdef _unroll #ifdef _unroll
#pragma unroll #pragma unroll
@ -139,7 +139,7 @@ u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4
return j; return j;
} }
int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], u32 timestamp_ct[8]) int decrypt_and_check (SCR_TYPE RC4_KEY *rc4_key, u32 data[4], u32 timestamp_ct[8])
{ {
rc4_init_16 (rc4_key, data); rc4_init_16 (rc4_key, data);
@ -285,8 +285,20 @@ __kernel void m07500_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru
COPY_PW (pws[gid]); COPY_PW (pws[gid]);
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64]; __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]; u32 checksum[4];
checksum[0] = krb5pa_bufs[digests_offset].checksum[0]; checksum[0] = krb5pa_bufs[digests_offset].checksum[0];
@ -327,7 +339,7 @@ __kernel void m07500_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru
kerb_prepare (ctx.h, checksum, digest); kerb_prepare (ctx.h, checksum, digest);
if (decrypt_and_check (&rc4_keys[lid], digest, timestamp_ct) == 1) if (decrypt_and_check (rc4_key, digest, timestamp_ct) == 1)
{ {
if (atomic_inc (&hashes_shown[digests_offset]) == 0) if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{ {
@ -354,8 +366,20 @@ __kernel void m07500_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru
COPY_PW (pws[gid]); COPY_PW (pws[gid]);
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64]; __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]; u32 checksum[4];
checksum[0] = krb5pa_bufs[digests_offset].checksum[0]; checksum[0] = krb5pa_bufs[digests_offset].checksum[0];
@ -396,7 +420,7 @@ __kernel void m07500_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru
kerb_prepare (ctx.h, checksum, digest); kerb_prepare (ctx.h, checksum, digest);
if (decrypt_and_check (&rc4_keys[lid], digest, timestamp_ct) == 1) if (decrypt_and_check (rc4_key, digest, timestamp_ct) == 1)
{ {
if (atomic_inc (&hashes_shown[digests_offset]) == 0) if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{ {

38
OpenCL/m07500_a1-optimized.cl
View File

@ -23,7 +23,7 @@ typedef struct
} RC4_KEY; } 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; u8 tmp;
@ -32,12 +32,12 @@ void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
rc4_key->S[j] = tmp; 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 v = 0x03020100;
u32 a = 0x04040404; u32 a = 0x04040404;
__local u32 *ptr = (__local u32 *) rc4_key->S; SCR_TYPE u32 *ptr = (SCR_TYPE u32 *) rc4_key->S;
#ifdef _unroll #ifdef _unroll
#pragma 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, const u32 in[4], u32 out[4]) u8 rc4_next_16 (SCR_TYPE RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4])
{ {
#ifdef _unroll #ifdef _unroll
#pragma unroll #pragma unroll
@ -138,7 +138,7 @@ u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4
return j; return j;
} }
int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], u32 timestamp_ct[8]) int decrypt_and_check (SCR_TYPE RC4_KEY *rc4_key, u32 data[4], u32 timestamp_ct[8])
{ {
rc4_init_16 (rc4_key, data); rc4_init_16 (rc4_key, data);
@ -438,8 +438,20 @@ __kernel void m07500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* shared * shared
*/ */
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64]; __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
/** /**
* loop * loop
*/ */
@ -518,7 +530,7 @@ __kernel void m07500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
tmp[2] = digest[2]; tmp[2] = digest[2];
tmp[3] = digest[3]; tmp[3] = digest[3];
if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1) if (decrypt_and_check (rc4_key, tmp, timestamp_ct) == 1)
{ {
if (atomic_inc (&hashes_shown[digests_offset]) == 0) if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{ {
@ -592,8 +604,20 @@ __kernel void m07500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* shared * shared
*/ */
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64]; __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
/** /**
* loop * loop
*/ */
@ -672,7 +696,7 @@ __kernel void m07500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
tmp[2] = digest[2]; tmp[2] = digest[2];
tmp[3] = digest[3]; tmp[3] = digest[3];
if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1) if (decrypt_and_check (rc4_key, tmp, timestamp_ct) == 1)
{ {
if (atomic_inc (&hashes_shown[digests_offset]) == 0) if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{ {

38
OpenCL/m07500_a1.cl
View File

@ -22,7 +22,7 @@ typedef struct
} RC4_KEY; } 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; u8 tmp;
@ -31,12 +31,12 @@ void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
rc4_key->S[j] = tmp; 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 v = 0x03020100;
u32 a = 0x04040404; u32 a = 0x04040404;
__local u32 *ptr = (__local u32 *) rc4_key->S; SCR_TYPE u32 *ptr = (SCR_TYPE u32 *) rc4_key->S;
#ifdef _unroll #ifdef _unroll
#pragma 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, const u32 in[4], u32 out[4]) u8 rc4_next_16 (SCR_TYPE RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4])
{ {
#ifdef _unroll #ifdef _unroll
#pragma unroll #pragma unroll
@ -137,7 +137,7 @@ u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4
return j; return j;
} }
int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], u32 timestamp_ct[8]) int decrypt_and_check (SCR_TYPE RC4_KEY *rc4_key, u32 data[4], u32 timestamp_ct[8])
{ {
rc4_init_16 (rc4_key, data); rc4_init_16 (rc4_key, data);
@ -281,8 +281,20 @@ __kernel void m07500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
* base * base
*/ */
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64]; __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]; u32 checksum[4];
checksum[0] = krb5pa_bufs[digests_offset].checksum[0]; checksum[0] = krb5pa_bufs[digests_offset].checksum[0];
@ -323,7 +335,7 @@ __kernel void m07500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
kerb_prepare (ctx.h, checksum, digest); kerb_prepare (ctx.h, checksum, digest);
if (decrypt_and_check (&rc4_keys[lid], digest, timestamp_ct) == 1) if (decrypt_and_check (rc4_key, digest, timestamp_ct) == 1)
{ {
if (atomic_inc (&hashes_shown[digests_offset]) == 0) if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{ {
@ -348,8 +360,20 @@ __kernel void m07500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
* base * base
*/ */
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64]; __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]; u32 checksum[4];
checksum[0] = krb5pa_bufs[digests_offset].checksum[0]; checksum[0] = krb5pa_bufs[digests_offset].checksum[0];
@ -390,7 +414,7 @@ __kernel void m07500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
kerb_prepare (ctx.h, checksum, digest); kerb_prepare (ctx.h, checksum, digest);
if (decrypt_and_check (&rc4_keys[lid], digest, timestamp_ct) == 1) if (decrypt_and_check (rc4_key, digest, timestamp_ct) == 1)
{ {
if (atomic_inc (&hashes_shown[digests_offset]) == 0) if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{ {

70
OpenCL/m07500_a3-optimized.cl
View File

@ -23,7 +23,7 @@ typedef struct
} RC4_KEY; } 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; u8 tmp;
@ -32,12 +32,12 @@ void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
rc4_key->S[j] = tmp; 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 v = 0x03020100;
u32 a = 0x04040404; u32 a = 0x04040404;
__local u32 *ptr = (__local u32 *) rc4_key->S; SCR_TYPE u32 *ptr = (SCR_TYPE u32 *) rc4_key->S;
#ifdef _unroll #ifdef _unroll
#pragma 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, const u32 in[4], u32 out[4]) u8 rc4_next_16 (SCR_TYPE RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4])
{ {
#ifdef _unroll #ifdef _unroll
#pragma unroll #pragma unroll
@ -138,7 +138,7 @@ u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4
return j; return j;
} }
int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], u32 timestamp_ct[8]) int decrypt_and_check (SCR_TYPE RC4_KEY *rc4_key, u32 data[4], u32 timestamp_ct[8])
{ {
rc4_init_16 (rc4_key, data); rc4_init_16 (rc4_key, data);
@ -382,7 +382,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); hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
} }
void m07500 (__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 krb5pa_t *krb5pa_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 m07500 (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 krb5pa_t *krb5pa_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 * modifier
@ -462,7 +462,7 @@ void m07500 (__local RC4_KEY *rc4_keys, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[
tmp[2] = digest[2]; tmp[2] = digest[2];
tmp[3] = digest[3]; tmp[3] = digest[3];
if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1) if (decrypt_and_check (rc4_key, tmp, timestamp_ct) == 1)
{ {
if (atomic_inc (&hashes_shown[digests_offset]) == 0) if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{ {
@ -517,9 +517,21 @@ __kernel void m07500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* main * main
*/ */
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64]; __local RC4_KEY rc4_keys[64];
m07500 (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, krb5pa_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
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, krb5pa_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 m07500_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 krb5pa_t *krb5pa_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) __kernel void m07500_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 krb5pa_t *krb5pa_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)
@ -567,9 +579,21 @@ __kernel void m07500_m08 (__global pw_t *pws, __global const kernel_rule_t *rule
* main * main
*/ */
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64]; __local RC4_KEY rc4_keys[64];
m07500 (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, krb5pa_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
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, krb5pa_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 m07500_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 krb5pa_t *krb5pa_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) __kernel void m07500_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 krb5pa_t *krb5pa_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)
@ -621,9 +645,21 @@ __kernel void m07500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* main * main
*/ */
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64]; __local RC4_KEY rc4_keys[64];
m07500 (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, krb5pa_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
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, krb5pa_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 m07500_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 krb5pa_t *krb5pa_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) __kernel void m07500_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 krb5pa_t *krb5pa_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)
@ -671,9 +707,21 @@ __kernel void m07500_s08 (__global pw_t *pws, __global const kernel_rule_t *rule
* main * main
*/ */
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64]; __local RC4_KEY rc4_keys[64];
m07500 (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, krb5pa_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
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, krb5pa_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 m07500_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 krb5pa_t *krb5pa_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) __kernel void m07500_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 krb5pa_t *krb5pa_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/m07500_a3.cl
View File

@ -22,7 +22,7 @@ typedef struct
} RC4_KEY; } 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; u8 tmp;
@ -31,12 +31,12 @@ void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
rc4_key->S[j] = tmp; 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 v = 0x03020100;
u32 a = 0x04040404; u32 a = 0x04040404;
__local u32 *ptr = (__local u32 *) rc4_key->S; SCR_TYPE u32 *ptr = (SCR_TYPE u32 *) rc4_key->S;
#ifdef _unroll #ifdef _unroll
#pragma 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, const u32 in[4], u32 out[4]) u8 rc4_next_16 (SCR_TYPE RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4])
{ {
#ifdef _unroll #ifdef _unroll
#pragma unroll #pragma unroll
@ -137,7 +137,7 @@ u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4
return j; return j;
} }
int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], u32 timestamp_ct[8]) int decrypt_and_check (SCR_TYPE RC4_KEY *rc4_key, u32 data[4], u32 timestamp_ct[8])
{ {
rc4_init_16 (rc4_key, data); rc4_init_16 (rc4_key, data);
@ -302,8 +302,20 @@ __kernel void m07500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
} }
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64]; __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]; u32 checksum[4];
checksum[0] = krb5pa_bufs[digests_offset].checksum[0]; checksum[0] = krb5pa_bufs[digests_offset].checksum[0];
@ -348,7 +360,7 @@ __kernel void m07500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
kerb_prepare (ctx.h, checksum, digest); kerb_prepare (ctx.h, checksum, digest);
if (decrypt_and_check (&rc4_keys[lid], digest, timestamp_ct) == 1) if (decrypt_and_check (rc4_key, digest, timestamp_ct) == 1)
{ {
if (atomic_inc (&hashes_shown[digests_offset]) == 0) if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{ {
@ -394,8 +406,20 @@ __kernel void m07500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
} }
#ifdef REAL_SHM
__local RC4_KEY rc4_keys[64]; __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]; u32 checksum[4];
checksum[0] = krb5pa_bufs[digests_offset].checksum[0]; checksum[0] = krb5pa_bufs[digests_offset].checksum[0];
@ -440,7 +464,7 @@ __kernel void m07500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
kerb_prepare (ctx.h, checksum, digest); kerb_prepare (ctx.h, checksum, digest);
if (decrypt_and_check (&rc4_keys[lid], digest, timestamp_ct) == 1) if (decrypt_and_check (rc4_key, digest, timestamp_ct) == 1)
{ {
if (atomic_inc (&hashes_shown[digests_offset]) == 0) if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{ {