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

Make -m 14000 a pure kernel only and add volatile for asm statement

Browse Source
This commit is contained in:
jsteube 2017-08-28 22:20:40 +02:00
parent 99f416435e
commit 2b9888486e
5 changed files with 143 additions and 191 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
OpenCL/m01500_a3.cl
View File

@ -55,7 +55,7 @@
// The underlying mathematical formulas are NOT copyrighted. // The underlying mathematical formulas are NOT copyrighted.
// //
#define LUT(a,b,c,d,e) u32 a; asm ("lop3.b32 %0, %1, %2, %3, "#e";" : "=r"(a): "r"(b), "r"(c), "r"(d)); #define LUT(a,b,c,d,e) u32 a; asm volatile ("lop3.b32 %0, %1, %2, %3, "#e";" : "=r"(a): "r"(b), "r"(c), "r"(d));
void s1 (const u32 a1, const u32 a2, const u32 a3, const u32 a4, const u32 a5, const u32 a6, u32 *out1, u32 *out2, u32 *out3, u32 *out4) void s1 (const u32 a1, const u32 a2, const u32 a3, const u32 a4, const u32 a5, const u32 a6, u32 *out1, u32 *out2, u32 *out3, u32 *out4)
{ {

2
OpenCL/m03000_a3.cl
View File

@ -52,7 +52,7 @@
// The underlying mathematical formulas are NOT copyrighted. // The underlying mathematical formulas are NOT copyrighted.
// //
#define LUT(a,b,c,d,e) u32 a; asm ("lop3.b32 %0, %1, %2, %3, "#e";" : "=r"(a): "r"(b), "r"(c), "r"(d)); #define LUT(a,b,c,d,e) u32 a; asm volatile ("lop3.b32 %0, %1, %2, %3, "#e";" : "=r"(a): "r"(b), "r"(c), "r"(d));
void s1 (const u32 a1, const u32 a2, const u32 a3, const u32 a4, const u32 a5, const u32 a6, u32 *out1, u32 *out2, u32 *out3, u32 *out4) void s1 (const u32 a1, const u32 a2, const u32 a3, const u32 a4, const u32 a5, const u32 a6, u32 *out1, u32 *out2, u32 *out3, u32 *out4)
{ {

140
OpenCL/m14000_a0-optimized.cl → OpenCL/m14000_a0.cl
View File

@ -3,7 +3,7 @@
* License.....: MIT * License.....: MIT
*/ */
#define NEW_SIMD_CODE //#define NEW_SIMD_CODE
#include "inc_vendor.cl" #include "inc_vendor.cl"
#include "inc_hash_constants.h" #include "inc_hash_constants.h"
@ -380,23 +380,23 @@ __constant u32a c_skb[8][64] =
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7], (S)[(i).s8], (S)[(i).s9], (S)[(i).sa], (S)[(i).sb], (S)[(i).sc], (S)[(i).sd], (S)[(i).se], (S)[(i).sf]) #define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7], (S)[(i).s8], (S)[(i).s9], (S)[(i).sa], (S)[(i).sb], (S)[(i).sc], (S)[(i).sd], (S)[(i).se], (S)[(i).sf])
#endif #endif
void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64]) void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64])
{ {
u32x r = rotl32 (data[0], 3u); u32 r = rotl32_S (data[0], 3u);
u32x l = rotl32 (data[1], 3u); u32 l = rotl32_S (data[1], 3u);
u32x tt; u32 tt;
#ifdef _unroll #ifdef _unroll
#pragma unroll #pragma unroll
#endif #endif
for (u32 i = 0; i < 16; i += 2) for (u32 i = 0; i < 16; i += 2)
{ {
u32x u; u32 u;
u32x t; u32 t;
u = Kc[i + 0] ^ r; u = Kc[i + 0] ^ r;
t = Kd[i + 0] ^ rotl32 (r, 28u); t = Kd[i + 0] ^ rotl32_S (r, 28u);
l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans) | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
@ -408,7 +408,7 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans); | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
u = Kc[i + 1] ^ l; u = Kc[i + 1] ^ l;
t = Kd[i + 1] ^ rotl32 (l, 28u); t = Kd[i + 1] ^ rotl32_S (l, 28u);
r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans) | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
@ -420,13 +420,13 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans); | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
} }
iv[0] = rotl32 (l, 29u); iv[0] = rotl32_S (l, 29u);
iv[1] = rotl32 (r, 29u); iv[1] = rotl32_S (r, 29u);
} }
void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64]) void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 (*s_skb)[64])
{ {
u32x tt; u32 tt;
PERM_OP (d, c, tt, 4, 0x0f0f0f0f); PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
HPERM_OP (c, tt, 2, 0xcccc0000); HPERM_OP (c, tt, 2, 0xcccc0000);
@ -461,42 +461,42 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
c = c & 0x0fffffff; c = c & 0x0fffffff;
d = d & 0x0fffffff; d = d & 0x0fffffff;
const u32x c00 = (c >> 0) & 0x0000003f; const u32 c00 = (c >> 0) & 0x0000003f;
const u32x c06 = (c >> 6) & 0x00383003; const u32 c06 = (c >> 6) & 0x00383003;
const u32x c07 = (c >> 7) & 0x0000003c; const u32 c07 = (c >> 7) & 0x0000003c;
const u32x c13 = (c >> 13) & 0x0000060f; const u32 c13 = (c >> 13) & 0x0000060f;
const u32x c20 = (c >> 20) & 0x00000001; const u32 c20 = (c >> 20) & 0x00000001;
u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb) u32 s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
| BOX (((c06 >> 0) & 0xff) | BOX (((c06 >> 0) & 0xff)
|((c07 >> 0) & 0xff), 1, s_skb) |((c07 >> 0) & 0xff), 1, s_skb)
| BOX (((c13 >> 0) & 0xff) | BOX (((c13 >> 0) & 0xff)
|((c06 >> 8) & 0xff), 2, s_skb) |((c06 >> 8) & 0xff), 2, s_skb)
| BOX (((c20 >> 0) & 0xff) | BOX (((c20 >> 0) & 0xff)
|((c13 >> 8) & 0xff) |((c13 >> 8) & 0xff)
|((c06 >> 16) & 0xff), 3, s_skb); |((c06 >> 16) & 0xff), 3, s_skb);
const u32x d00 = (d >> 0) & 0x00003c3f; const u32 d00 = (d >> 0) & 0x00003c3f;
const u32x d07 = (d >> 7) & 0x00003f03; const u32 d07 = (d >> 7) & 0x00003f03;
const u32x d21 = (d >> 21) & 0x0000000f; const u32 d21 = (d >> 21) & 0x0000000f;
const u32x d22 = (d >> 22) & 0x00000030; const u32 d22 = (d >> 22) & 0x00000030;
u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb) u32 t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
| BOX (((d07 >> 0) & 0xff) | BOX (((d07 >> 0) & 0xff)
|((d00 >> 8) & 0xff), 5, s_skb) |((d00 >> 8) & 0xff), 5, s_skb)
| BOX (((d07 >> 8) & 0xff), 6, s_skb) | BOX (((d07 >> 8) & 0xff), 6, s_skb)
| BOX (((d21 >> 0) & 0xff) | BOX (((d21 >> 0) & 0xff)
|((d22 >> 0) & 0xff), 7, s_skb); |((d22 >> 0) & 0xff), 7, s_skb);
Kc[i] = ((t << 16) | (s & 0x0000ffff)); Kc[i] = ((t << 16) | (s & 0x0000ffff));
Kd[i] = ((s >> 16) | (t & 0xffff0000)); Kd[i] = ((s >> 16) | (t & 0xffff0000));
Kc[i] = rotl32 (Kc[i], 2u); Kc[i] = rotl32_S (Kc[i], 2u);
Kd[i] = rotl32 (Kd[i], 2u); Kd[i] = rotl32_S (Kd[i], 2u);
} }
} }
__kernel void m14000_m04 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) __kernel void m14000_mxx (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{ {
/** /**
* base * base
@ -571,45 +571,37 @@ __kernel void m14000_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
{ {
u32x w0[4] = { 0 }; u32 w0[4] = { 0 };
u32x w1[4] = { 0 }; u32 w1[4] = { 0 };
u32x w2[4] = { 0 }; u32 w2[4] = { 0 };
u32x w3[4] = { 0 }; u32 w3[4] = { 0 };
apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
const u32x c = w0[0]; const u32 c = w0[0];
const u32x d = w0[1]; const u32 d = w0[1];
u32x Kc[16]; u32 Kc[16];
u32x Kd[16]; u32 Kd[16];
_des_crypt_keysetup (c, d, Kc, Kd, s_skb); _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
u32x data[2]; u32 data[2];
data[0] = salt_buf0[0]; data[0] = salt_buf0[0];
data[1] = salt_buf0[1]; data[1] = salt_buf0[1];
u32x iv[2]; u32 iv[2];
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
u32x z = 0; u32 z = 0;
COMPARE_M_SIMD (iv[0], iv[1], z, z); COMPARE_M_SIMD (iv[0], iv[1], z, z);
} }
} }
__kernel void m14000_m08 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) __kernel void m14000_sxx (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{
}
__kernel void m14000_m16 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{
}
__kernel void m14000_s04 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{ {
/** /**
* base * base
@ -696,40 +688,32 @@ __kernel void m14000_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
{ {
u32x w0[4] = { 0 }; u32 w0[4] = { 0 };
u32x w1[4] = { 0 }; u32 w1[4] = { 0 };
u32x w2[4] = { 0 }; u32 w2[4] = { 0 };
u32x w3[4] = { 0 }; u32 w3[4] = { 0 };
apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
const u32x c = w0[0]; const u32 c = w0[0];
const u32x d = w0[1]; const u32 d = w0[1];
u32x Kc[16]; u32 Kc[16];
u32x Kd[16]; u32 Kd[16];
_des_crypt_keysetup (c, d, Kc, Kd, s_skb); _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
u32x data[2]; u32 data[2];
data[0] = salt_buf0[0]; data[0] = salt_buf0[0];
data[1] = salt_buf0[1]; data[1] = salt_buf0[1];
u32x iv[2]; u32 iv[2];
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
u32x z = 0; u32 z = 0;
COMPARE_S_SIMD (iv[0], iv[1], z, z); COMPARE_S_SIMD (iv[0], iv[1], z, z);
} }
} }
__kernel void m14000_s08 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{
}
__kernel void m14000_s16 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{
}

168
OpenCL/m14000_a1-optimized.cl → OpenCL/m14000_a1.cl
View File

@ -3,7 +3,7 @@
* License.....: MIT * License.....: MIT
*/ */
#define NEW_SIMD_CODE //#define NEW_SIMD_CODE
#include "inc_vendor.cl" #include "inc_vendor.cl"
#include "inc_hash_constants.h" #include "inc_hash_constants.h"
@ -370,23 +370,23 @@ __constant u32a c_skb[8][64] =
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7], (S)[(i).s8], (S)[(i).s9], (S)[(i).sa], (S)[(i).sb], (S)[(i).sc], (S)[(i).sd], (S)[(i).se], (S)[(i).sf]) #define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7], (S)[(i).s8], (S)[(i).s9], (S)[(i).sa], (S)[(i).sb], (S)[(i).sc], (S)[(i).sd], (S)[(i).se], (S)[(i).sf])
#endif #endif
void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64]) void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64])
{ {
u32x r = rotl32 (data[0], 3u); u32 r = rotl32_S (data[0], 3u);
u32x l = rotl32 (data[1], 3u); u32 l = rotl32_S (data[1], 3u);
u32x tt; u32 tt;
#ifdef _unroll #ifdef _unroll
#pragma unroll #pragma unroll
#endif #endif
for (u32 i = 0; i < 16; i += 2) for (u32 i = 0; i < 16; i += 2)
{ {
u32x u; u32 u;
u32x t; u32 t;
u = Kc[i + 0] ^ r; u = Kc[i + 0] ^ r;
t = Kd[i + 0] ^ rotl32 (r, 28u); t = Kd[i + 0] ^ rotl32_S (r, 28u);
l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans) | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
@ -398,7 +398,7 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans); | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
u = Kc[i + 1] ^ l; u = Kc[i + 1] ^ l;
t = Kd[i + 1] ^ rotl32 (l, 28u); t = Kd[i + 1] ^ rotl32_S (l, 28u);
r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans) | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
@ -410,13 +410,13 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans); | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
} }
iv[0] = rotl32 (l, 29u); iv[0] = rotl32_S (l, 29u);
iv[1] = rotl32 (r, 29u); iv[1] = rotl32_S (r, 29u);
} }
void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64]) void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 (*s_skb)[64])
{ {
u32x tt; u32 tt;
PERM_OP (d, c, tt, 4, 0x0f0f0f0f); PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
HPERM_OP (c, tt, 2, 0xcccc0000); HPERM_OP (c, tt, 2, 0xcccc0000);
@ -451,42 +451,42 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
c = c & 0x0fffffff; c = c & 0x0fffffff;
d = d & 0x0fffffff; d = d & 0x0fffffff;
const u32x c00 = (c >> 0) & 0x0000003f; const u32 c00 = (c >> 0) & 0x0000003f;
const u32x c06 = (c >> 6) & 0x00383003; const u32 c06 = (c >> 6) & 0x00383003;
const u32x c07 = (c >> 7) & 0x0000003c; const u32 c07 = (c >> 7) & 0x0000003c;
const u32x c13 = (c >> 13) & 0x0000060f; const u32 c13 = (c >> 13) & 0x0000060f;
const u32x c20 = (c >> 20) & 0x00000001; const u32 c20 = (c >> 20) & 0x00000001;
u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb) u32 s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
| BOX (((c06 >> 0) & 0xff) | BOX (((c06 >> 0) & 0xff)
|((c07 >> 0) & 0xff), 1, s_skb) |((c07 >> 0) & 0xff), 1, s_skb)
| BOX (((c13 >> 0) & 0xff) | BOX (((c13 >> 0) & 0xff)
|((c06 >> 8) & 0xff), 2, s_skb) |((c06 >> 8) & 0xff), 2, s_skb)
| BOX (((c20 >> 0) & 0xff) | BOX (((c20 >> 0) & 0xff)
|((c13 >> 8) & 0xff) |((c13 >> 8) & 0xff)
|((c06 >> 16) & 0xff), 3, s_skb); |((c06 >> 16) & 0xff), 3, s_skb);
const u32x d00 = (d >> 0) & 0x00003c3f; const u32 d00 = (d >> 0) & 0x00003c3f;
const u32x d07 = (d >> 7) & 0x00003f03; const u32 d07 = (d >> 7) & 0x00003f03;
const u32x d21 = (d >> 21) & 0x0000000f; const u32 d21 = (d >> 21) & 0x0000000f;
const u32x d22 = (d >> 22) & 0x00000030; const u32 d22 = (d >> 22) & 0x00000030;
u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb) u32 t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
| BOX (((d07 >> 0) & 0xff) | BOX (((d07 >> 0) & 0xff)
|((d00 >> 8) & 0xff), 5, s_skb) |((d00 >> 8) & 0xff), 5, s_skb)
| BOX (((d07 >> 8) & 0xff), 6, s_skb) | BOX (((d07 >> 8) & 0xff), 6, s_skb)
| BOX (((d21 >> 0) & 0xff) | BOX (((d21 >> 0) & 0xff)
|((d22 >> 0) & 0xff), 7, s_skb); |((d22 >> 0) & 0xff), 7, s_skb);
Kc[i] = ((t << 16) | (s & 0x0000ffff)); Kc[i] = ((t << 16) | (s & 0x0000ffff));
Kd[i] = ((s >> 16) | (t & 0xffff0000)); Kd[i] = ((s >> 16) | (t & 0xffff0000));
Kc[i] = rotl32 (Kc[i], 2u); Kc[i] = rotl32_S (Kc[i], 2u);
Kd[i] = rotl32 (Kd[i], 2u); Kd[i] = rotl32_S (Kd[i], 2u);
} }
} }
__kernel void m14000_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) __kernel void m14000_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{ {
/** /**
* base * base
@ -561,18 +561,18 @@ __kernel void m14000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
{ {
const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
const u32x pw_len = pw_l_len + pw_r_len; const u32 pw_len = pw_l_len + pw_r_len;
/** /**
* concat password candidate * concat password candidate
*/ */
u32x wordl0[4] = { 0 }; u32 wordl0[4] = { 0 };
u32x wordl1[4] = { 0 }; u32 wordl1[4] = { 0 };
u32x wordl2[4] = { 0 }; u32 wordl2[4] = { 0 };
u32x wordl3[4] = { 0 }; u32 wordl3[4] = { 0 };
wordl0[0] = pw_buf0[0]; wordl0[0] = pw_buf0[0];
wordl0[1] = pw_buf0[1]; wordl0[1] = pw_buf0[1];
@ -583,10 +583,10 @@ __kernel void m14000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
wordl1[2] = pw_buf1[2]; wordl1[2] = pw_buf1[2];
wordl1[3] = pw_buf1[3]; wordl1[3] = pw_buf1[3];
u32x wordr0[4] = { 0 }; u32 wordr0[4] = { 0 };
u32x wordr1[4] = { 0 }; u32 wordr1[4] = { 0 };
u32x wordr2[4] = { 0 }; u32 wordr2[4] = { 0 };
u32x wordr3[4] = { 0 }; u32 wordr3[4] = { 0 };
wordr0[0] = ix_create_combt (combs_buf, il_pos, 0); wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
wordr0[1] = ix_create_combt (combs_buf, il_pos, 1); wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
@ -606,43 +606,35 @@ __kernel void m14000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len); switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
} }
u32x w0[2]; u32 w0[2];
w0[0] = wordl0[0] | wordr0[0]; w0[0] = wordl0[0] | wordr0[0];
w0[1] = wordl0[1] | wordr0[1]; w0[1] = wordl0[1] | wordr0[1];
const u32x c = w0[0]; const u32 c = w0[0];
const u32x d = w0[1]; const u32 d = w0[1];
u32x Kc[16]; u32 Kc[16];
u32x Kd[16]; u32 Kd[16];
_des_crypt_keysetup (c, d, Kc, Kd, s_skb); _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
u32x data[2]; u32 data[2];
data[0] = salt_buf0[0]; data[0] = salt_buf0[0];
data[1] = salt_buf0[1]; data[1] = salt_buf0[1];
u32x iv[2]; u32 iv[2];
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
u32x z = 0; u32 z = 0;
COMPARE_M_SIMD (iv[0], iv[1], z, z); COMPARE_M_SIMD (iv[0], iv[1], z, z);
} }
} }
__kernel void m14000_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 const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) __kernel void m14000_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{
}
__kernel void m14000_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 const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{
}
__kernel void m14000_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{ {
/** /**
* base * base
@ -729,18 +721,18 @@ __kernel void m14000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
{ {
const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
const u32x pw_len = pw_l_len + pw_r_len; const u32 pw_len = pw_l_len + pw_r_len;
/** /**
* concat password candidate * concat password candidate
*/ */
u32x wordl0[4] = { 0 }; u32 wordl0[4] = { 0 };
u32x wordl1[4] = { 0 }; u32 wordl1[4] = { 0 };
u32x wordl2[4] = { 0 }; u32 wordl2[4] = { 0 };
u32x wordl3[4] = { 0 }; u32 wordl3[4] = { 0 };
wordl0[0] = pw_buf0[0]; wordl0[0] = pw_buf0[0];
wordl0[1] = pw_buf0[1]; wordl0[1] = pw_buf0[1];
@ -751,10 +743,10 @@ __kernel void m14000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
wordl1[2] = pw_buf1[2]; wordl1[2] = pw_buf1[2];
wordl1[3] = pw_buf1[3]; wordl1[3] = pw_buf1[3];
u32x wordr0[4] = { 0 }; u32 wordr0[4] = { 0 };
u32x wordr1[4] = { 0 }; u32 wordr1[4] = { 0 };
u32x wordr2[4] = { 0 }; u32 wordr2[4] = { 0 };
u32x wordr3[4] = { 0 }; u32 wordr3[4] = { 0 };
wordr0[0] = ix_create_combt (combs_buf, il_pos, 0); wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
wordr0[1] = ix_create_combt (combs_buf, il_pos, 1); wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
@ -774,38 +766,30 @@ __kernel void m14000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len); switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
} }
u32x w0[2]; u32 w0[2];
w0[0] = wordl0[0] | wordr0[0]; w0[0] = wordl0[0] | wordr0[0];
w0[1] = wordl0[1] | wordr0[1]; w0[1] = wordl0[1] | wordr0[1];
const u32x c = (w0[0]); const u32 c = (w0[0]);
const u32x d = (w0[1]); const u32 d = (w0[1]);
u32x Kc[16]; u32 Kc[16];
u32x Kd[16]; u32 Kd[16];
_des_crypt_keysetup (c, d, Kc, Kd, s_skb); _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
u32x data[2]; u32 data[2];
data[0] = salt_buf0[0]; data[0] = salt_buf0[0];
data[1] = salt_buf0[1]; data[1] = salt_buf0[1];
u32x iv[2]; u32 iv[2];
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
u32x z = 0; u32 z = 0;
COMPARE_S_SIMD (iv[0], iv[1], z, z); COMPARE_S_SIMD (iv[0], iv[1], z, z);
} }
} }
__kernel void m14000_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 const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{
}
__kernel void m14000_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 const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{
}

22
OpenCL/m14000_a3-optimized.cl → OpenCL/m14000_a3.cl
View File

@ -52,7 +52,7 @@
// The underlying mathematical formulas are NOT copyrighted. // The underlying mathematical formulas are NOT copyrighted.
// //
#define LUT(a,b,c,d,e) u32 a; asm ("lop3.b32 %0, %1, %2, %3, "#e";" : "=r"(a): "r"(b), "r"(c), "r"(d)); #define LUT(a,b,c,d,e) u32 a; asm volatile ("lop3.b32 %0, %1, %2, %3, "#e";" : "=r"(a): "r"(b), "r"(c), "r"(d));
void s1 (const u32 a1, const u32 a2, const u32 a3, const u32 a4, const u32 a5, const u32 a6, u32 *out1, u32 *out2, u32 *out3, u32 *out4) void s1 (const u32 a1, const u32 a2, const u32 a3, const u32 a4, const u32 a5, const u32 a6, u32 *out1, u32 *out2, u32 *out3, u32 *out4)
{ {
@ -2714,7 +2714,7 @@ __kernel void m14000_tm (__global u32 *mod, __global bs_word_t *words_buf_r)
} }
} }
__kernel void m14000_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bs_word_t *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) __kernel void m14000_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bs_word_t *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{ {
/** /**
* base * base
@ -2732,15 +2732,7 @@ __kernel void m14000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
m14000m (pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); m14000m (pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
} }
__kernel void m14000_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bs_word_t *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) __kernel void m14000_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bs_word_t *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{
}
__kernel void m14000_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bs_word_t *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{
}
__kernel void m14000_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bs_word_t *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{ {
/** /**
* base * base
@ -2757,11 +2749,3 @@ __kernel void m14000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
m14000s (pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); m14000s (pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
} }
__kernel void m14000_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bs_word_t *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{
}
__kernel void m14000_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bs_word_t *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{
}