diff --git a/OpenCL/m13100_a0-optimized.cl b/OpenCL/m13100_a0-optimized.cl index f1a566a8e..3d0e8bb36 100644 --- a/OpenCL/m13100_a0-optimized.cl +++ b/OpenCL/m13100_a0-optimized.cl @@ -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) { diff --git a/OpenCL/m13100_a0.cl b/OpenCL/m13100_a0.cl index 270cef3b5..a82bbdc49 100644 --- a/OpenCL/m13100_a0.cl +++ b/OpenCL/m13100_a0.cl @@ -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) { diff --git a/OpenCL/m13100_a1-optimized.cl b/OpenCL/m13100_a1-optimized.cl index 7f864abaa..f2314427d 100644 --- a/OpenCL/m13100_a1-optimized.cl +++ b/OpenCL/m13100_a1-optimized.cl @@ -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) { diff --git a/OpenCL/m13100_a1.cl b/OpenCL/m13100_a1.cl index 2d0c113b5..dc3075af3 100644 --- a/OpenCL/m13100_a1.cl +++ b/OpenCL/m13100_a1.cl @@ -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) { diff --git a/OpenCL/m13100_a3-optimized.cl b/OpenCL/m13100_a3-optimized.cl index ab8e2825c..b397fa378 100644 --- a/OpenCL/m13100_a3-optimized.cl +++ b/OpenCL/m13100_a3-optimized.cl @@ -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) diff --git a/OpenCL/m13100_a3.cl b/OpenCL/m13100_a3.cl index 14d78b0c2..0db38f6ba 100644 --- a/OpenCL/m13100_a3.cl +++ b/OpenCL/m13100_a3.cl @@ -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) {