#include "OpenCL/common.c"
#include "OpenCL/simd.c"
+#define ROUND(v) \
+{ \
+ a ^= (((a & 0x3f) + add) * (v)) + (a << 8); \
+ b += (b << 8) ^ a; \
+ add += v; \
+}
+
+#define CODE_PRE \
+{ \
+ for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) \
+ { \
+ const u32x w0r = words_buf_r[il_pos / VECT_SIZE]; \
+ \
+ const u32x w0 = w0l | w0r; \
+ \
+ u32x a = MYSQL323_A; \
+ u32x b = MYSQL323_B; \
+ u32x c = 0; \
+ u32x d = 0; \
+ \
+ u32x add = 7; \
+
+#define CODE_LOOP(rest) \
+ \
+ int i; \
+ int j; \
+ \
+ for (i = 0, j = 1; i <= (int) (rest) - 4; i += 4, j += 1) \
+ { \
+ const u32 wj = w[j]; \
+ \
+ ROUND ((wj >> 0) & 0xff); \
+ ROUND ((wj >> 8) & 0xff); \
+ ROUND ((wj >> 16) & 0xff); \
+ ROUND ((wj >> 24) & 0xff); \
+ } \
+ \
+ const u32 wj = w[j]; \
+ \
+ const u32 left = (rest) - i; \
+ \
+ if (left == 3) \
+ { \
+ ROUND ((wj >> 0) & 0xff); \
+ ROUND ((wj >> 8) & 0xff); \
+ ROUND ((wj >> 16) & 0xff); \
+ } \
+ else if (left == 2) \
+ { \
+ ROUND ((wj >> 0) & 0xff); \
+ ROUND ((wj >> 8) & 0xff); \
+ } \
+ else if (left == 1) \
+ { \
+ ROUND ((wj >> 0) & 0xff); \
+ }
+
+#define CODE_POST_M \
+ \
+ a &= 0x7fffffff; \
+ b &= 0x7fffffff; \
+ \
+ COMPARE_M_SIMD (a, b, c, d); \
+ } \
+}
+
+#define CODE_POST_S \
+ \
+ a &= 0x7fffffff; \
+ b &= 0x7fffffff; \
+ \
+ COMPARE_S_SIMD (a, b, c, d); \
+ } \
+}
+
static void m00200m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 lid = get_local_id (0);
/**
- * loop
+ * digest
*/
- u32 w0l = w[0];
-
- for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
+ const u32 search[4] =
{
- const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
-
- const u32x w0 = w0l | w0r;
-
- u32x w_t[16];
-
- w_t[ 0] = w0;
- w_t[ 1] = w[ 1];
- w_t[ 2] = w[ 2];
- w_t[ 3] = w[ 3];
- w_t[ 4] = w[ 4];
- w_t[ 5] = w[ 5];
- w_t[ 6] = w[ 6];
- w_t[ 7] = w[ 7];
- w_t[ 8] = w[ 8];
- w_t[ 9] = w[ 9];
- w_t[10] = w[10];
- w_t[11] = w[11];
- w_t[12] = w[12];
- w_t[13] = w[13];
- w_t[14] = w[14];
- w_t[15] = w[15];
-
- u32x a = MYSQL323_A;
- u32x b = MYSQL323_B;
- u32x c = 0;
- u32x d = 0;
-
- u32x add = 7;
-
- #define ROUND(v) \
- { \
- a ^= (((a & 0x3f) + add) * (v)) + (a << 8); \
- b += (b << 8) ^ a; \
- add += v; \
- }
-
- int i;
- int j;
-
- for (i = 0, j = 0; i <= (int) pw_len - 4; i += 4, j += 1)
- {
- const u32x wj = w_t[j];
-
- ROUND ((wj >> 0) & 0xff);
- ROUND ((wj >> 8) & 0xff);
- ROUND ((wj >> 16) & 0xff);
- ROUND ((wj >> 24) & 0xff);
- }
-
- const u32x wj = w_t[j];
-
- const u32 left = pw_len - i;
+ digests_buf[digests_offset].digest_buf[DGST_R0],
+ digests_buf[digests_offset].digest_buf[DGST_R1],
+ digests_buf[digests_offset].digest_buf[DGST_R2],
+ digests_buf[digests_offset].digest_buf[DGST_R3]
+ };
- if (left == 3)
- {
- ROUND ((wj >> 0) & 0xff);
- ROUND ((wj >> 8) & 0xff);
- ROUND ((wj >> 16) & 0xff);
- }
- else if (left == 2)
- {
- ROUND ((wj >> 0) & 0xff);
- ROUND ((wj >> 8) & 0xff);
- }
- else if (left == 1)
- {
- ROUND ((wj >> 0) & 0xff);
- }
+ /**
+ * loop
+ */
- a &= 0x7fffffff;
- b &= 0x7fffffff;
+ u32 w0l = w[0];
- COMPARE_M_SIMD (a, b, c, d);
+ switch (pw_len)
+ {
+ case 1:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff);
+ CODE_POST_M;
+ break;
+
+ case 2:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff);
+ CODE_POST_M;
+ break;
+
+ case 3:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff);
+ CODE_POST_M;
+ break;
+
+ case 4:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ CODE_POST_M;
+ break;
+
+ case 5:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff);
+ CODE_POST_M;
+ break;
+
+ case 6:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff);
+ CODE_POST_M;
+ break;
+
+ case 7:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff);
+ CODE_POST_M;
+ break;
+
+ case 8:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ CODE_POST_M;
+ break;
+
+ case 9:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ ROUND ((w[2] >> 0) & 0xff);
+ CODE_POST_M;
+ break;
+
+ case 10:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ ROUND ((w[2] >> 0) & 0xff); ROUND ((w[2] >> 8) & 0xff);
+ CODE_POST_M;
+ break;
+
+ case 11:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ ROUND ((w[2] >> 0) & 0xff); ROUND ((w[2] >> 8) & 0xff); ROUND ((w[2] >> 16) & 0xff);
+ CODE_POST_M;
+ break;
+
+ case 12:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ ROUND ((w[2] >> 0) & 0xff); ROUND ((w[2] >> 8) & 0xff); ROUND ((w[2] >> 16) & 0xff); ROUND ((w[2] >> 24) & 0xff);
+ CODE_POST_M;
+ break;
+
+ case 13:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ ROUND ((w[2] >> 0) & 0xff); ROUND ((w[2] >> 8) & 0xff); ROUND ((w[2] >> 16) & 0xff); ROUND ((w[2] >> 24) & 0xff);
+ ROUND ((w[3] >> 0) & 0xff);
+ CODE_POST_M;
+ break;
+
+ case 14:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ ROUND ((w[2] >> 0) & 0xff); ROUND ((w[2] >> 8) & 0xff); ROUND ((w[2] >> 16) & 0xff); ROUND ((w[2] >> 24) & 0xff);
+ ROUND ((w[3] >> 0) & 0xff); ROUND ((w[3] >> 8) & 0xff);
+ CODE_POST_M;
+ break;
+
+ case 15:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ ROUND ((w[2] >> 0) & 0xff); ROUND ((w[2] >> 8) & 0xff); ROUND ((w[2] >> 16) & 0xff); ROUND ((w[2] >> 24) & 0xff);
+ ROUND ((w[3] >> 0) & 0xff); ROUND ((w[3] >> 8) & 0xff); ROUND ((w[3] >> 16) & 0xff);
+ CODE_POST_M;
+ break;
+
+ case 16:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ ROUND ((w[2] >> 0) & 0xff); ROUND ((w[2] >> 8) & 0xff); ROUND ((w[2] >> 16) & 0xff); ROUND ((w[2] >> 24) & 0xff);
+ ROUND ((w[3] >> 0) & 0xff); ROUND ((w[3] >> 8) & 0xff); ROUND ((w[3] >> 16) & 0xff); ROUND ((w[3] >> 24) & 0xff);
+ CODE_POST_M;
+ break;
+
+ default:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ CODE_LOOP (pw_len - 4);
+ CODE_POST_M;
+ break;
}
}
u32 w0l = w[0];
- for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
+ switch (pw_len)
{
- const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
-
- const u32x w0 = w0l | w0r;
-
- u32x w_t[16];
-
- w_t[ 0] = w0;
- w_t[ 1] = w[ 1];
- w_t[ 2] = w[ 2];
- w_t[ 3] = w[ 3];
- w_t[ 4] = w[ 4];
- w_t[ 5] = w[ 5];
- w_t[ 6] = w[ 6];
- w_t[ 7] = w[ 7];
- w_t[ 8] = w[ 8];
- w_t[ 9] = w[ 9];
- w_t[10] = w[10];
- w_t[11] = w[11];
- w_t[12] = w[12];
- w_t[13] = w[13];
- w_t[14] = w[14];
- w_t[15] = w[15];
-
- u32x a = MYSQL323_A;
- u32x b = MYSQL323_B;
- u32x c = 0;
- u32x d = 0;
-
- u32x add = 7;
-
- #define ROUND(v) \
- { \
- a ^= (((a & 0x3f) + add) * (v)) + (a << 8); \
- b += (b << 8) ^ a; \
- add += v; \
- }
-
- int i;
- int j;
-
- for (i = 0, j = 0; i <= (int) pw_len - 4; i += 4, j += 1)
- {
- const u32x wj = w_t[j];
-
- ROUND ((wj >> 0) & 0xff);
- ROUND ((wj >> 8) & 0xff);
- ROUND ((wj >> 16) & 0xff);
- ROUND ((wj >> 24) & 0xff);
- }
-
- const u32x wj = w_t[j];
-
- const u32 left = pw_len - i;
-
- if (left == 3)
- {
- ROUND ((wj >> 0) & 0xff);
- ROUND ((wj >> 8) & 0xff);
- ROUND ((wj >> 16) & 0xff);
- }
- else if (left == 2)
- {
- ROUND ((wj >> 0) & 0xff);
- ROUND ((wj >> 8) & 0xff);
- }
- else if (left == 1)
- {
- ROUND ((wj >> 0) & 0xff);
- }
-
- a &= 0x7fffffff;
- b &= 0x7fffffff;
-
- COMPARE_S_SIMD (a, b, c, d);
+ case 1:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff);
+ CODE_POST_S;
+ break;
+
+ case 2:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff);
+ CODE_POST_S;
+ break;
+
+ case 3:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff);
+ CODE_POST_S;
+ break;
+
+ case 4:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ CODE_POST_S;
+ break;
+
+ case 5:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff);
+ CODE_POST_S;
+ break;
+
+ case 6:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff);
+ CODE_POST_S;
+ break;
+
+ case 7:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff);
+ CODE_POST_S;
+ break;
+
+ case 8:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ CODE_POST_S;
+ break;
+
+ case 9:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ ROUND ((w[2] >> 0) & 0xff);
+ CODE_POST_S;
+ break;
+
+ case 10:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ ROUND ((w[2] >> 0) & 0xff); ROUND ((w[2] >> 8) & 0xff);
+ CODE_POST_S;
+ break;
+
+ case 11:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ ROUND ((w[2] >> 0) & 0xff); ROUND ((w[2] >> 8) & 0xff); ROUND ((w[2] >> 16) & 0xff);
+ CODE_POST_S;
+ break;
+
+ case 12:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ ROUND ((w[2] >> 0) & 0xff); ROUND ((w[2] >> 8) & 0xff); ROUND ((w[2] >> 16) & 0xff); ROUND ((w[2] >> 24) & 0xff);
+ CODE_POST_S;
+ break;
+
+ case 13:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ ROUND ((w[2] >> 0) & 0xff); ROUND ((w[2] >> 8) & 0xff); ROUND ((w[2] >> 16) & 0xff); ROUND ((w[2] >> 24) & 0xff);
+ ROUND ((w[3] >> 0) & 0xff);
+ CODE_POST_S;
+ break;
+
+ case 14:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ ROUND ((w[2] >> 0) & 0xff); ROUND ((w[2] >> 8) & 0xff); ROUND ((w[2] >> 16) & 0xff); ROUND ((w[2] >> 24) & 0xff);
+ ROUND ((w[3] >> 0) & 0xff); ROUND ((w[3] >> 8) & 0xff);
+ CODE_POST_S;
+ break;
+
+ case 15:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ ROUND ((w[2] >> 0) & 0xff); ROUND ((w[2] >> 8) & 0xff); ROUND ((w[2] >> 16) & 0xff); ROUND ((w[2] >> 24) & 0xff);
+ ROUND ((w[3] >> 0) & 0xff); ROUND ((w[3] >> 8) & 0xff); ROUND ((w[3] >> 16) & 0xff);
+ CODE_POST_S;
+ break;
+
+ case 16:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ ROUND ((w[1] >> 0) & 0xff); ROUND ((w[1] >> 8) & 0xff); ROUND ((w[1] >> 16) & 0xff); ROUND ((w[1] >> 24) & 0xff);
+ ROUND ((w[2] >> 0) & 0xff); ROUND ((w[2] >> 8) & 0xff); ROUND ((w[2] >> 16) & 0xff); ROUND ((w[2] >> 24) & 0xff);
+ ROUND ((w[3] >> 0) & 0xff); ROUND ((w[3] >> 8) & 0xff); ROUND ((w[3] >> 16) & 0xff); ROUND ((w[3] >> 24) & 0xff);
+ CODE_POST_S;
+ break;
+
+ default:
+ CODE_PRE;
+ ROUND ((w0 >> 0) & 0xff); ROUND ((w0 >> 8) & 0xff); ROUND ((w0 >> 16) & 0xff); ROUND ((w0 >> 24) & 0xff);
+ CODE_LOOP (pw_len - 4);
+ CODE_POST_S;
+ break;
}
}
#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));
-static 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)
{
LUT(xAA55AA5500550055, a1, a4, a6, 0xC1)
LUT(xA55AA55AF0F5F0F5, a3, a6, xAA55AA5500550055, 0x9E)
*out4 ^= x4;
}
-static void s2 (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 s2 (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)
{
LUT(xEEEEEEEE99999999, a1, a2, a6, 0x97)
LUT(xFFFFEEEE66666666, a5, a6, xEEEEEEEE99999999, 0x67)
*out4 ^= x4;
}
-static void s3 (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 s3 (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)
{
LUT(xA50FA50FA50FA50F, a1, a3, a4, 0xC9)
LUT(xF0F00F0FF0F0F0F0, a3, a5, a6, 0x4B)
*out4 ^= x4;
}
-static void s4 (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 s4 (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)
{
LUT(x55F055F055F055F0, a1, a3, a4, 0x72)
LUT(xA500F5F0A500F5F0, a3, a5, x55F055F055F055F0, 0xAD)
*out4 ^= x4;
}
-static void s5 (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 s5 (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)
{
LUT(xA0A0A0A0FFFFFFFF, a1, a3, a6, 0xAB)
LUT(xFFFF00005555FFFF, a1, a5, a6, 0xB9)
*out4 ^= x4;
}
-static void s6 (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 s6 (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)
{
LUT(x5050F5F55050F5F5, a1, a3, a5, 0xB2)
LUT(x6363C6C66363C6C6, a1, a2, x5050F5F55050F5F5, 0x66)
*out4 ^= x4;
}
-static void s7 (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 s7 (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)
{
LUT(x88AA88AA88AA88AA, a1, a2, a4, 0x0B)
LUT(xAAAAFF00AAAAFF00, a1, a4, a5, 0x27)
*out4 ^= x4;
}
-static void s8 (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 s8 (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)
{
LUT(xEEEE3333EEEE3333, a1, a2, a5, 0x9D)
LUT(xBBBBBBBBBBBBBBBB, a1, a1, a2, 0x83)
* The effort has been sponsored by Rapid7: http://www.rapid7.com
*/
-static 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)
{
u32 x55005500, x5A0F5A0F, x3333FFFF, x66666666, x22226666, x2D2D6969,
x25202160;
*out4 ^= x31;
}
-static void s2 (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 s2 (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)
{
u32 x33CC33CC;
u32 x55550000, x00AA00FF, x33BB33FF;
*out4 ^= x31;
}
-static void s3 (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 s3 (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)
{
u32 x44444444, x0F0FF0F0, x4F4FF4F4, x00FFFF00, x00AAAA00, x4FE55EF4;
u32 x3C3CC3C3, x3C3C0000, x7373F4F4, x0C840A00;
*out3 ^= x21;
}
-static void s4 (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 s4 (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)
{
u32 x5A5A5A5A, x0F0FF0F0;
u32 x33FF33FF, x33FFCC00, x0C0030F0, x0C0CC0C0, x0CF3C03F, x5EFBDA7F,
*out4 ^= x31;
}
-static void s5 (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 s5 (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)
{
u32 x77777777, x77770000, x22225555, x11116666, x1F1F6F6F;
u32 x70700000, x43433333, x00430033, x55557777, x55167744, x5A19784B;
*out2 ^= x11;
}
-static void s6 (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 s6 (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)
{
u32 x33CC33CC;
u32 x3333FFFF, x11115555, x22DD6699, x22DD9966, x00220099;
*out1 ^= x01;
}
-static void s7 (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 s7 (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)
{
u32 x0FF00FF0, x3CC33CC3, x00003CC3, x0F000F00, x5A555A55, x00001841;
u32 x00000F00, x33333C33, x7B777E77, x0FF0F00F, x74878E78;
*out2 ^= x11;
}
-static void s8 (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 s8 (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)
{
u32 x0C0C0C0C, x0000F0F0, x00FFF00F, x00555005, x00515001;
u32 x33000330, x77555775, x30303030, x3030CFCF, x30104745, x30555745;
* The effort has been sponsored by Rapid7: http://www.rapid7.com
*/
-static 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)
{
u32 x55005500, x5A0F5A0F, x3333FFFF, x66666666, x22226666, x2D2D6969,
x25202160;
*out4 ^= x31;
}
-static void s2 (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 s2 (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)
{
u32 x33CC33CC;
u32 x55550000, x00AA00FF, x33BB33FF;
*out4 ^= x31;
}
-static void s3 (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 s3 (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)
{
u32 x44444444, x0F0FF0F0, x4F4FF4F4, x00FFFF00, x00AAAA00, x4FE55EF4;
u32 x3C3CC3C3, x3C3C0000, x7373F4F4, x0C840A00;
*out3 ^= x21;
}
-static void s4 (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 s4 (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)
{
u32 x5A5A5A5A, x0F0FF0F0;
u32 x33FF33FF, x33FFCC00, x0C0030F0, x0C0CC0C0, x0CF3C03F, x5EFBDA7F,
*out4 ^= x31;
}
-static void s5 (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 s5 (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)
{
u32 x77777777, x77770000, x22225555, x11116666, x1F1F6F6F;
u32 x70700000, x43433333, x00430033, x55557777, x55167744, x5A19784B;
*out2 ^= x11;
}
-static void s6 (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 s6 (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)
{
u32 x33CC33CC;
u32 x3333FFFF, x11115555, x22DD6699, x22DD9966, x00220099;
*out1 ^= x01;
}
-static void s7 (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 s7 (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)
{
u32 x0FF00FF0, x3CC33CC3, x00003CC3, x0F000F00, x5A555A55, x00001841;
u32 x00000F00, x33333C33, x7B777E77, x0FF0F00F, x74878E78;
*out2 ^= x11;
}
-static void s8 (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 s8 (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)
{
u32 x0C0C0C0C, x0000F0F0, x00FFF00F, x00555005, x00515001;
u32 x33000330, x77555775, x30303030, x3030CFCF, x30104745, x30555745;
#ifdef DESCRYPT_SALT
-static void DESCrypt (const u32 SALT, const u32 K00, const u32 K01, const u32 K02, const u32 K03, const u32 K04, const u32 K05, const u32 K06, const u32 K07, const u32 K08, const u32 K09, const u32 K10, const u32 K11, const u32 K12, const u32 K13, const u32 K14, const u32 K15, const u32 K16, const u32 K17, const u32 K18, const u32 K19, const u32 K20, const u32 K21, const u32 K22, const u32 K23, const u32 K24, const u32 K25, const u32 K26, const u32 K27, const u32 K28, const u32 K29, const u32 K30, const u32 K31, const u32 K32, const u32 K33, const u32 K34, const u32 K35, const u32 K36, const u32 K37, const u32 K38, const u32 K39, const u32 K40, const u32 K41, const u32 K42, const u32 K43, const u32 K44, const u32 K45, const u32 K46, const u32 K47, const u32 K48, const u32 K49, const u32 K50, const u32 K51, const u32 K52, const u32 K53, const u32 K54, const u32 K55, u32 *D00, u32 *D01, u32 *D02, u32 *D03, u32 *D04, u32 *D05, u32 *D06, u32 *D07, u32 *D08, u32 *D09, u32 *D10, u32 *D11, u32 *D12, u32 *D13, u32 *D14, u32 *D15, u32 *D16, u32 *D17, u32 *D18, u32 *D19, u32 *D20, u32 *D21, u32 *D22, u32 *D23, u32 *D24, u32 *D25, u32 *D26, u32 *D27, u32 *D28, u32 *D29, u32 *D30, u32 *D31, u32 *D32, u32 *D33, u32 *D34, u32 *D35, u32 *D36, u32 *D37, u32 *D38, u32 *D39, u32 *D40, u32 *D41, u32 *D42, u32 *D43, u32 *D44, u32 *D45, u32 *D46, u32 *D47, u32 *D48, u32 *D49, u32 *D50, u32 *D51, u32 *D52, u32 *D53, u32 *D54, u32 *D55, u32 *D56, u32 *D57, u32 *D58, u32 *D59, u32 *D60, u32 *D61, u32 *D62, u32 *D63)
+void DESCrypt (const u32 SALT, const u32 K00, const u32 K01, const u32 K02, const u32 K03, const u32 K04, const u32 K05, const u32 K06, const u32 K07, const u32 K08, const u32 K09, const u32 K10, const u32 K11, const u32 K12, const u32 K13, const u32 K14, const u32 K15, const u32 K16, const u32 K17, const u32 K18, const u32 K19, const u32 K20, const u32 K21, const u32 K22, const u32 K23, const u32 K24, const u32 K25, const u32 K26, const u32 K27, const u32 K28, const u32 K29, const u32 K30, const u32 K31, const u32 K32, const u32 K33, const u32 K34, const u32 K35, const u32 K36, const u32 K37, const u32 K38, const u32 K39, const u32 K40, const u32 K41, const u32 K42, const u32 K43, const u32 K44, const u32 K45, const u32 K46, const u32 K47, const u32 K48, const u32 K49, const u32 K50, const u32 K51, const u32 K52, const u32 K53, const u32 K54, const u32 K55, u32 *D00, u32 *D01, u32 *D02, u32 *D03, u32 *D04, u32 *D05, u32 *D06, u32 *D07, u32 *D08, u32 *D09, u32 *D10, u32 *D11, u32 *D12, u32 *D13, u32 *D14, u32 *D15, u32 *D16, u32 *D17, u32 *D18, u32 *D19, u32 *D20, u32 *D21, u32 *D22, u32 *D23, u32 *D24, u32 *D25, u32 *D26, u32 *D27, u32 *D28, u32 *D29, u32 *D30, u32 *D31, u32 *D32, u32 *D33, u32 *D34, u32 *D35, u32 *D36, u32 *D37, u32 *D38, u32 *D39, u32 *D40, u32 *D41, u32 *D42, u32 *D43, u32 *D44, u32 *D45, u32 *D46, u32 *D47, u32 *D48, u32 *D49, u32 *D50, u32 *D51, u32 *D52, u32 *D53, u32 *D54, u32 *D55, u32 *D56, u32 *D57, u32 *D58, u32 *D59, u32 *D60, u32 *D61, u32 *D62, u32 *D63)
{
sXXX_DECL u32 s001 = (0x001 & DESCRYPT_SALT) ? 0xffffffff : 0;
sXXX_DECL u32 s002 = (0x002 & DESCRYPT_SALT) ? 0xffffffff : 0;
#else
-static void DESCrypt (const u32 SALT, const u32 K00, const u32 K01, const u32 K02, const u32 K03, const u32 K04, const u32 K05, const u32 K06, const u32 K07, const u32 K08, const u32 K09, const u32 K10, const u32 K11, const u32 K12, const u32 K13, const u32 K14, const u32 K15, const u32 K16, const u32 K17, const u32 K18, const u32 K19, const u32 K20, const u32 K21, const u32 K22, const u32 K23, const u32 K24, const u32 K25, const u32 K26, const u32 K27, const u32 K28, const u32 K29, const u32 K30, const u32 K31, const u32 K32, const u32 K33, const u32 K34, const u32 K35, const u32 K36, const u32 K37, const u32 K38, const u32 K39, const u32 K40, const u32 K41, const u32 K42, const u32 K43, const u32 K44, const u32 K45, const u32 K46, const u32 K47, const u32 K48, const u32 K49, const u32 K50, const u32 K51, const u32 K52, const u32 K53, const u32 K54, const u32 K55, u32 *D00, u32 *D01, u32 *D02, u32 *D03, u32 *D04, u32 *D05, u32 *D06, u32 *D07, u32 *D08, u32 *D09, u32 *D10, u32 *D11, u32 *D12, u32 *D13, u32 *D14, u32 *D15, u32 *D16, u32 *D17, u32 *D18, u32 *D19, u32 *D20, u32 *D21, u32 *D22, u32 *D23, u32 *D24, u32 *D25, u32 *D26, u32 *D27, u32 *D28, u32 *D29, u32 *D30, u32 *D31, u32 *D32, u32 *D33, u32 *D34, u32 *D35, u32 *D36, u32 *D37, u32 *D38, u32 *D39, u32 *D40, u32 *D41, u32 *D42, u32 *D43, u32 *D44, u32 *D45, u32 *D46, u32 *D47, u32 *D48, u32 *D49, u32 *D50, u32 *D51, u32 *D52, u32 *D53, u32 *D54, u32 *D55, u32 *D56, u32 *D57, u32 *D58, u32 *D59, u32 *D60, u32 *D61, u32 *D62, u32 *D63)
+void DESCrypt (const u32 SALT, const u32 K00, const u32 K01, const u32 K02, const u32 K03, const u32 K04, const u32 K05, const u32 K06, const u32 K07, const u32 K08, const u32 K09, const u32 K10, const u32 K11, const u32 K12, const u32 K13, const u32 K14, const u32 K15, const u32 K16, const u32 K17, const u32 K18, const u32 K19, const u32 K20, const u32 K21, const u32 K22, const u32 K23, const u32 K24, const u32 K25, const u32 K26, const u32 K27, const u32 K28, const u32 K29, const u32 K30, const u32 K31, const u32 K32, const u32 K33, const u32 K34, const u32 K35, const u32 K36, const u32 K37, const u32 K38, const u32 K39, const u32 K40, const u32 K41, const u32 K42, const u32 K43, const u32 K44, const u32 K45, const u32 K46, const u32 K47, const u32 K48, const u32 K49, const u32 K50, const u32 K51, const u32 K52, const u32 K53, const u32 K54, const u32 K55, u32 *D00, u32 *D01, u32 *D02, u32 *D03, u32 *D04, u32 *D05, u32 *D06, u32 *D07, u32 *D08, u32 *D09, u32 *D10, u32 *D11, u32 *D12, u32 *D13, u32 *D14, u32 *D15, u32 *D16, u32 *D17, u32 *D18, u32 *D19, u32 *D20, u32 *D21, u32 *D22, u32 *D23, u32 *D24, u32 *D25, u32 *D26, u32 *D27, u32 *D28, u32 *D29, u32 *D30, u32 *D31, u32 *D32, u32 *D33, u32 *D34, u32 *D35, u32 *D36, u32 *D37, u32 *D38, u32 *D39, u32 *D40, u32 *D41, u32 *D42, u32 *D43, u32 *D44, u32 *D45, u32 *D46, u32 *D47, u32 *D48, u32 *D49, u32 *D50, u32 *D51, u32 *D52, u32 *D53, u32 *D54, u32 *D55, u32 *D56, u32 *D57, u32 *D58, u32 *D59, u32 *D60, u32 *D61, u32 *D62, u32 *D63)
{
sXXX_DECL u32 s001 = (0x001 & SALT) ? 0xffffffff : 0;
sXXX_DECL u32 s002 = (0x002 & SALT) ? 0xffffffff : 0;
#endif
-static void transpose32c (u32 data[32])
+void transpose32c (u32 data[32])
{
#define swap(x,y,j,m) \
t = ((x) ^ ((y) >> (j))) & (m); \
swap (data[30], data[31], 1, 0x55555555);
}
-static void m01500m (__local u32 *s_S, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 m01500m (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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)
{
/**
* base
const u32 w0s = (w0 << 1) & 0xfefefefe;
const u32 w1s = (w1 << 1) & 0xfefefefe;
- const u32 K00 = -((w0s >> ( 0 + 7)) & 1);
- const u32 K01 = -((w0s >> ( 0 + 6)) & 1);
- const u32 K02 = -((w0s >> ( 0 + 5)) & 1);
- const u32 K03 = -((w0s >> ( 0 + 4)) & 1);
- const u32 K04 = -((w0s >> ( 0 + 3)) & 1);
- const u32 K05 = -((w0s >> ( 0 + 2)) & 1);
- const u32 K06 = -((w0s >> ( 0 + 1)) & 1);
- const u32 K07 = -((w0s >> ( 8 + 7)) & 1);
- const u32 K08 = -((w0s >> ( 8 + 6)) & 1);
- const u32 K09 = -((w0s >> ( 8 + 5)) & 1);
- const u32 K10 = -((w0s >> ( 8 + 4)) & 1);
- const u32 K11 = -((w0s >> ( 8 + 3)) & 1);
- const u32 K12 = -((w0s >> ( 8 + 2)) & 1);
- const u32 K13 = -((w0s >> ( 8 + 1)) & 1);
- const u32 K14 = -((w0s >> (16 + 7)) & 1);
- const u32 K15 = -((w0s >> (16 + 6)) & 1);
- const u32 K16 = -((w0s >> (16 + 5)) & 1);
- const u32 K17 = -((w0s >> (16 + 4)) & 1);
- const u32 K18 = -((w0s >> (16 + 3)) & 1);
- const u32 K19 = -((w0s >> (16 + 2)) & 1);
- const u32 K20 = -((w0s >> (16 + 1)) & 1);
- const u32 K21 = -((w0s >> (24 + 7)) & 1);
- const u32 K22 = -((w0s >> (24 + 6)) & 1);
- const u32 K23 = -((w0s >> (24 + 5)) & 1);
- const u32 K24 = -((w0s >> (24 + 4)) & 1);
- const u32 K25 = -((w0s >> (24 + 3)) & 1);
- const u32 K26 = -((w0s >> (24 + 2)) & 1);
- const u32 K27 = -((w0s >> (24 + 1)) & 1);
- const u32 K28 = -((w1s >> ( 0 + 7)) & 1);
- const u32 K29 = -((w1s >> ( 0 + 6)) & 1);
- const u32 K30 = -((w1s >> ( 0 + 5)) & 1);
- const u32 K31 = -((w1s >> ( 0 + 4)) & 1);
- const u32 K32 = -((w1s >> ( 0 + 3)) & 1);
- const u32 K33 = -((w1s >> ( 0 + 2)) & 1);
- const u32 K34 = -((w1s >> ( 0 + 1)) & 1);
- const u32 K35 = -((w1s >> ( 8 + 7)) & 1);
- const u32 K36 = -((w1s >> ( 8 + 6)) & 1);
- const u32 K37 = -((w1s >> ( 8 + 5)) & 1);
- const u32 K38 = -((w1s >> ( 8 + 4)) & 1);
- const u32 K39 = -((w1s >> ( 8 + 3)) & 1);
- const u32 K40 = -((w1s >> ( 8 + 2)) & 1);
- const u32 K41 = -((w1s >> ( 8 + 1)) & 1);
- const u32 K42 = -((w1s >> (16 + 7)) & 1);
- const u32 K43 = -((w1s >> (16 + 6)) & 1);
- const u32 K44 = -((w1s >> (16 + 5)) & 1);
- const u32 K45 = -((w1s >> (16 + 4)) & 1);
- const u32 K46 = -((w1s >> (16 + 3)) & 1);
- const u32 K47 = -((w1s >> (16 + 2)) & 1);
- const u32 K48 = -((w1s >> (16 + 1)) & 1);
- const u32 K49 = -((w1s >> (24 + 7)) & 1);
- const u32 K50 = -((w1s >> (24 + 6)) & 1);
- const u32 K51 = -((w1s >> (24 + 5)) & 1);
- const u32 K52 = -((w1s >> (24 + 4)) & 1);
- const u32 K53 = -((w1s >> (24 + 3)) & 1);
- const u32 K54 = -((w1s >> (24 + 2)) & 1);
- const u32 K55 = -((w1s >> (24 + 1)) & 1);
+ #define K00 -((w0s >> ( 0 + 7)) & 1)
+ #define K01 -((w0s >> ( 0 + 6)) & 1)
+ #define K02 -((w0s >> ( 0 + 5)) & 1)
+ #define K03 -((w0s >> ( 0 + 4)) & 1)
+ #define K04 -((w0s >> ( 0 + 3)) & 1)
+ #define K05 -((w0s >> ( 0 + 2)) & 1)
+ #define K06 -((w0s >> ( 0 + 1)) & 1)
+ #define K07 -((w0s >> ( 8 + 7)) & 1)
+ #define K08 -((w0s >> ( 8 + 6)) & 1)
+ #define K09 -((w0s >> ( 8 + 5)) & 1)
+ #define K10 -((w0s >> ( 8 + 4)) & 1)
+ #define K11 -((w0s >> ( 8 + 3)) & 1)
+ #define K12 -((w0s >> ( 8 + 2)) & 1)
+ #define K13 -((w0s >> ( 8 + 1)) & 1)
+ #define K14 -((w0s >> (16 + 7)) & 1)
+ #define K15 -((w0s >> (16 + 6)) & 1)
+ #define K16 -((w0s >> (16 + 5)) & 1)
+ #define K17 -((w0s >> (16 + 4)) & 1)
+ #define K18 -((w0s >> (16 + 3)) & 1)
+ #define K19 -((w0s >> (16 + 2)) & 1)
+ #define K20 -((w0s >> (16 + 1)) & 1)
+ #define K21 -((w0s >> (24 + 7)) & 1)
+ #define K22 -((w0s >> (24 + 6)) & 1)
+ #define K23 -((w0s >> (24 + 5)) & 1)
+ #define K24 -((w0s >> (24 + 4)) & 1)
+ #define K25 -((w0s >> (24 + 3)) & 1)
+ #define K26 -((w0s >> (24 + 2)) & 1)
+ #define K27 -((w0s >> (24 + 1)) & 1)
+ #define K28 -((w1s >> ( 0 + 7)) & 1)
+ #define K29 -((w1s >> ( 0 + 6)) & 1)
+ #define K30 -((w1s >> ( 0 + 5)) & 1)
+ #define K31 -((w1s >> ( 0 + 4)) & 1)
+ #define K32 -((w1s >> ( 0 + 3)) & 1)
+ #define K33 -((w1s >> ( 0 + 2)) & 1)
+ #define K34 -((w1s >> ( 0 + 1)) & 1)
+ #define K35 -((w1s >> ( 8 + 7)) & 1)
+ #define K36 -((w1s >> ( 8 + 6)) & 1)
+ #define K37 -((w1s >> ( 8 + 5)) & 1)
+ #define K38 -((w1s >> ( 8 + 4)) & 1)
+ #define K39 -((w1s >> ( 8 + 3)) & 1)
+ #define K40 -((w1s >> ( 8 + 2)) & 1)
+ #define K41 -((w1s >> ( 8 + 1)) & 1)
+ #define K42 -((w1s >> (16 + 7)) & 1)
+ #define K43 -((w1s >> (16 + 6)) & 1)
+ #define K44 -((w1s >> (16 + 5)) & 1)
+ #define K45 -((w1s >> (16 + 4)) & 1)
+ #define K46 -((w1s >> (16 + 3)) & 1)
+ #define K47 -((w1s >> (16 + 2)) & 1)
+ #define K48 -((w1s >> (16 + 1)) & 1)
+ #define K49 -((w1s >> (24 + 7)) & 1)
+ #define K50 -((w1s >> (24 + 6)) & 1)
+ #define K51 -((w1s >> (24 + 5)) & 1)
+ #define K52 -((w1s >> (24 + 4)) & 1)
+ #define K53 -((w1s >> (24 + 3)) & 1)
+ #define K54 -((w1s >> (24 + 2)) & 1)
+ #define K55 -((w1s >> (24 + 1)) & 1)
/**
* inner loop
}
}
-static void m01500s (__local u32 *s_S, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 m01500s (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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)
{
/**
* base
* digest
*/
- #define S00 s_S[ 0]
- #define S01 s_S[ 1]
- #define S02 s_S[ 2]
- #define S03 s_S[ 3]
- #define S04 s_S[ 4]
- #define S05 s_S[ 5]
- #define S06 s_S[ 6]
- #define S07 s_S[ 7]
- #define S08 s_S[ 8]
- #define S09 s_S[ 9]
- #define S10 s_S[10]
- #define S11 s_S[11]
- #define S12 s_S[12]
- #define S13 s_S[13]
- #define S14 s_S[14]
- #define S15 s_S[15]
- #define S16 s_S[16]
- #define S17 s_S[17]
- #define S18 s_S[18]
- #define S19 s_S[19]
- #define S20 s_S[20]
- #define S21 s_S[21]
- #define S22 s_S[22]
- #define S23 s_S[23]
- #define S24 s_S[24]
- #define S25 s_S[25]
- #define S26 s_S[26]
- #define S27 s_S[27]
- #define S28 s_S[28]
- #define S29 s_S[29]
- #define S30 s_S[30]
- #define S31 s_S[31]
- #define S32 s_S[32]
- #define S33 s_S[33]
- #define S34 s_S[34]
- #define S35 s_S[35]
- #define S36 s_S[36]
- #define S37 s_S[37]
- #define S38 s_S[38]
- #define S39 s_S[39]
- #define S40 s_S[40]
- #define S41 s_S[41]
- #define S42 s_S[42]
- #define S43 s_S[43]
- #define S44 s_S[44]
- #define S45 s_S[45]
- #define S46 s_S[46]
- #define S47 s_S[47]
- #define S48 s_S[48]
- #define S49 s_S[49]
- #define S50 s_S[50]
- #define S51 s_S[51]
- #define S52 s_S[52]
- #define S53 s_S[53]
- #define S54 s_S[54]
- #define S55 s_S[55]
- #define S56 s_S[56]
- #define S57 s_S[57]
- #define S58 s_S[58]
- #define S59 s_S[59]
- #define S60 s_S[60]
- #define S61 s_S[61]
- #define S62 s_S[62]
- #define S63 s_S[63]
+ const u32 s0 = digests_buf[0].digest_buf[0];
+ const u32 s1 = digests_buf[0].digest_buf[1];
+
+ #define S00 -((s0 >> 0) & 1)
+ #define S01 -((s0 >> 1) & 1)
+ #define S02 -((s0 >> 2) & 1)
+ #define S03 -((s0 >> 3) & 1)
+ #define S04 -((s0 >> 4) & 1)
+ #define S05 -((s0 >> 5) & 1)
+ #define S06 -((s0 >> 6) & 1)
+ #define S07 -((s0 >> 7) & 1)
+ #define S08 -((s0 >> 8) & 1)
+ #define S09 -((s0 >> 9) & 1)
+ #define S10 -((s0 >> 10) & 1)
+ #define S11 -((s0 >> 11) & 1)
+ #define S12 -((s0 >> 12) & 1)
+ #define S13 -((s0 >> 13) & 1)
+ #define S14 -((s0 >> 14) & 1)
+ #define S15 -((s0 >> 15) & 1)
+ #define S16 -((s0 >> 16) & 1)
+ #define S17 -((s0 >> 17) & 1)
+ #define S18 -((s0 >> 18) & 1)
+ #define S19 -((s0 >> 19) & 1)
+ #define S20 -((s0 >> 20) & 1)
+ #define S21 -((s0 >> 21) & 1)
+ #define S22 -((s0 >> 22) & 1)
+ #define S23 -((s0 >> 23) & 1)
+ #define S24 -((s0 >> 24) & 1)
+ #define S25 -((s0 >> 25) & 1)
+ #define S26 -((s0 >> 26) & 1)
+ #define S27 -((s0 >> 27) & 1)
+ #define S28 -((s0 >> 28) & 1)
+ #define S29 -((s0 >> 29) & 1)
+ #define S30 -((s0 >> 30) & 1)
+ #define S31 -((s0 >> 31) & 1)
+ #define S32 -((s1 >> 0) & 1)
+ #define S33 -((s1 >> 1) & 1)
+ #define S34 -((s1 >> 2) & 1)
+ #define S35 -((s1 >> 3) & 1)
+ #define S36 -((s1 >> 4) & 1)
+ #define S37 -((s1 >> 5) & 1)
+ #define S38 -((s1 >> 6) & 1)
+ #define S39 -((s1 >> 7) & 1)
+ #define S40 -((s1 >> 8) & 1)
+ #define S41 -((s1 >> 9) & 1)
+ #define S42 -((s1 >> 10) & 1)
+ #define S43 -((s1 >> 11) & 1)
+ #define S44 -((s1 >> 12) & 1)
+ #define S45 -((s1 >> 13) & 1)
+ #define S46 -((s1 >> 14) & 1)
+ #define S47 -((s1 >> 15) & 1)
+ #define S48 -((s1 >> 16) & 1)
+ #define S49 -((s1 >> 17) & 1)
+ #define S50 -((s1 >> 18) & 1)
+ #define S51 -((s1 >> 19) & 1)
+ #define S52 -((s1 >> 20) & 1)
+ #define S53 -((s1 >> 21) & 1)
+ #define S54 -((s1 >> 22) & 1)
+ #define S55 -((s1 >> 23) & 1)
+ #define S56 -((s1 >> 24) & 1)
+ #define S57 -((s1 >> 25) & 1)
+ #define S58 -((s1 >> 26) & 1)
+ #define S59 -((s1 >> 27) & 1)
+ #define S60 -((s1 >> 28) & 1)
+ #define S61 -((s1 >> 29) & 1)
+ #define S62 -((s1 >> 30) & 1)
+ #define S63 -((s1 >> 31) & 1)
/**
* base
const u32 w0s = (w0 << 1) & 0xfefefefe;
const u32 w1s = (w1 << 1) & 0xfefefefe;
- const u32 K00 = -((w0s >> ( 0 + 7)) & 1);
- const u32 K01 = -((w0s >> ( 0 + 6)) & 1);
- const u32 K02 = -((w0s >> ( 0 + 5)) & 1);
- const u32 K03 = -((w0s >> ( 0 + 4)) & 1);
- const u32 K04 = -((w0s >> ( 0 + 3)) & 1);
- const u32 K05 = -((w0s >> ( 0 + 2)) & 1);
- const u32 K06 = -((w0s >> ( 0 + 1)) & 1);
- const u32 K07 = -((w0s >> ( 8 + 7)) & 1);
- const u32 K08 = -((w0s >> ( 8 + 6)) & 1);
- const u32 K09 = -((w0s >> ( 8 + 5)) & 1);
- const u32 K10 = -((w0s >> ( 8 + 4)) & 1);
- const u32 K11 = -((w0s >> ( 8 + 3)) & 1);
- const u32 K12 = -((w0s >> ( 8 + 2)) & 1);
- const u32 K13 = -((w0s >> ( 8 + 1)) & 1);
- const u32 K14 = -((w0s >> (16 + 7)) & 1);
- const u32 K15 = -((w0s >> (16 + 6)) & 1);
- const u32 K16 = -((w0s >> (16 + 5)) & 1);
- const u32 K17 = -((w0s >> (16 + 4)) & 1);
- const u32 K18 = -((w0s >> (16 + 3)) & 1);
- const u32 K19 = -((w0s >> (16 + 2)) & 1);
- const u32 K20 = -((w0s >> (16 + 1)) & 1);
- const u32 K21 = -((w0s >> (24 + 7)) & 1);
- const u32 K22 = -((w0s >> (24 + 6)) & 1);
- const u32 K23 = -((w0s >> (24 + 5)) & 1);
- const u32 K24 = -((w0s >> (24 + 4)) & 1);
- const u32 K25 = -((w0s >> (24 + 3)) & 1);
- const u32 K26 = -((w0s >> (24 + 2)) & 1);
- const u32 K27 = -((w0s >> (24 + 1)) & 1);
- const u32 K28 = -((w1s >> ( 0 + 7)) & 1);
- const u32 K29 = -((w1s >> ( 0 + 6)) & 1);
- const u32 K30 = -((w1s >> ( 0 + 5)) & 1);
- const u32 K31 = -((w1s >> ( 0 + 4)) & 1);
- const u32 K32 = -((w1s >> ( 0 + 3)) & 1);
- const u32 K33 = -((w1s >> ( 0 + 2)) & 1);
- const u32 K34 = -((w1s >> ( 0 + 1)) & 1);
- const u32 K35 = -((w1s >> ( 8 + 7)) & 1);
- const u32 K36 = -((w1s >> ( 8 + 6)) & 1);
- const u32 K37 = -((w1s >> ( 8 + 5)) & 1);
- const u32 K38 = -((w1s >> ( 8 + 4)) & 1);
- const u32 K39 = -((w1s >> ( 8 + 3)) & 1);
- const u32 K40 = -((w1s >> ( 8 + 2)) & 1);
- const u32 K41 = -((w1s >> ( 8 + 1)) & 1);
- const u32 K42 = -((w1s >> (16 + 7)) & 1);
- const u32 K43 = -((w1s >> (16 + 6)) & 1);
- const u32 K44 = -((w1s >> (16 + 5)) & 1);
- const u32 K45 = -((w1s >> (16 + 4)) & 1);
- const u32 K46 = -((w1s >> (16 + 3)) & 1);
- const u32 K47 = -((w1s >> (16 + 2)) & 1);
- const u32 K48 = -((w1s >> (16 + 1)) & 1);
- const u32 K49 = -((w1s >> (24 + 7)) & 1);
- const u32 K50 = -((w1s >> (24 + 6)) & 1);
- const u32 K51 = -((w1s >> (24 + 5)) & 1);
- const u32 K52 = -((w1s >> (24 + 4)) & 1);
- const u32 K53 = -((w1s >> (24 + 3)) & 1);
- const u32 K54 = -((w1s >> (24 + 2)) & 1);
- const u32 K55 = -((w1s >> (24 + 1)) & 1);
+ #define K00 -((w0s >> ( 0 + 7)) & 1)
+ #define K01 -((w0s >> ( 0 + 6)) & 1)
+ #define K02 -((w0s >> ( 0 + 5)) & 1)
+ #define K03 -((w0s >> ( 0 + 4)) & 1)
+ #define K04 -((w0s >> ( 0 + 3)) & 1)
+ #define K05 -((w0s >> ( 0 + 2)) & 1)
+ #define K06 -((w0s >> ( 0 + 1)) & 1)
+ #define K07 -((w0s >> ( 8 + 7)) & 1)
+ #define K08 -((w0s >> ( 8 + 6)) & 1)
+ #define K09 -((w0s >> ( 8 + 5)) & 1)
+ #define K10 -((w0s >> ( 8 + 4)) & 1)
+ #define K11 -((w0s >> ( 8 + 3)) & 1)
+ #define K12 -((w0s >> ( 8 + 2)) & 1)
+ #define K13 -((w0s >> ( 8 + 1)) & 1)
+ #define K14 -((w0s >> (16 + 7)) & 1)
+ #define K15 -((w0s >> (16 + 6)) & 1)
+ #define K16 -((w0s >> (16 + 5)) & 1)
+ #define K17 -((w0s >> (16 + 4)) & 1)
+ #define K18 -((w0s >> (16 + 3)) & 1)
+ #define K19 -((w0s >> (16 + 2)) & 1)
+ #define K20 -((w0s >> (16 + 1)) & 1)
+ #define K21 -((w0s >> (24 + 7)) & 1)
+ #define K22 -((w0s >> (24 + 6)) & 1)
+ #define K23 -((w0s >> (24 + 5)) & 1)
+ #define K24 -((w0s >> (24 + 4)) & 1)
+ #define K25 -((w0s >> (24 + 3)) & 1)
+ #define K26 -((w0s >> (24 + 2)) & 1)
+ #define K27 -((w0s >> (24 + 1)) & 1)
+ #define K28 -((w1s >> ( 0 + 7)) & 1)
+ #define K29 -((w1s >> ( 0 + 6)) & 1)
+ #define K30 -((w1s >> ( 0 + 5)) & 1)
+ #define K31 -((w1s >> ( 0 + 4)) & 1)
+ #define K32 -((w1s >> ( 0 + 3)) & 1)
+ #define K33 -((w1s >> ( 0 + 2)) & 1)
+ #define K34 -((w1s >> ( 0 + 1)) & 1)
+ #define K35 -((w1s >> ( 8 + 7)) & 1)
+ #define K36 -((w1s >> ( 8 + 6)) & 1)
+ #define K37 -((w1s >> ( 8 + 5)) & 1)
+ #define K38 -((w1s >> ( 8 + 4)) & 1)
+ #define K39 -((w1s >> ( 8 + 3)) & 1)
+ #define K40 -((w1s >> ( 8 + 2)) & 1)
+ #define K41 -((w1s >> ( 8 + 1)) & 1)
+ #define K42 -((w1s >> (16 + 7)) & 1)
+ #define K43 -((w1s >> (16 + 6)) & 1)
+ #define K44 -((w1s >> (16 + 5)) & 1)
+ #define K45 -((w1s >> (16 + 4)) & 1)
+ #define K46 -((w1s >> (16 + 3)) & 1)
+ #define K47 -((w1s >> (16 + 2)) & 1)
+ #define K48 -((w1s >> (16 + 1)) & 1)
+ #define K49 -((w1s >> (24 + 7)) & 1)
+ #define K50 -((w1s >> (24 + 6)) & 1)
+ #define K51 -((w1s >> (24 + 5)) & 1)
+ #define K52 -((w1s >> (24 + 4)) & 1)
+ #define K53 -((w1s >> (24 + 3)) & 1)
+ #define K54 -((w1s >> (24 + 2)) & 1)
+ #define K55 -((w1s >> (24 + 1)) & 1)
/**
* inner loop
tmpResult |= D13 ^ S13;
tmpResult |= D14 ^ S14;
tmpResult |= D15 ^ S15;
+
+ if (tmpResult == 0xffffffff) return;
+
tmpResult |= D16 ^ S16;
tmpResult |= D17 ^ S17;
tmpResult |= D18 ^ S18;
tmpResult |= D29 ^ S29;
tmpResult |= D30 ^ S30;
tmpResult |= D31 ^ S31;
+
+ if (tmpResult == 0xffffffff) return;
+
tmpResult |= D32 ^ S32;
tmpResult |= D33 ^ S33;
tmpResult |= D34 ^ S34;
const u32 gid = get_global_id (0);
const u32 lid = get_local_id (0);
- const u32 vid = get_local_id (1);
-
- const u32 s0 = digests_buf[digests_offset].digest_buf[0];
- const u32 s1 = digests_buf[digests_offset].digest_buf[1];
-
- __local u32 s_S[64];
-
- if (lid == 0)
- {
- s_S[ 0 + vid] = -((s0 >> vid) & 1);
- }
- else if (lid == 1)
- {
- s_S[32 + vid] = -((s1 >> vid) & 1);
- }
-
- barrier (CLK_LOCAL_MEM_FENCE);
if (gid >= gid_max) return;
* main
*/
- m01500m (s_S, 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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
+ m01500m (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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
}
__kernel void m01500_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 u32 gid_max)
const u32 gid = get_global_id (0);
const u32 lid = get_local_id (0);
- const u32 vid = get_local_id (1);
-
- const u32 s0 = digests_buf[digests_offset].digest_buf[0];
- const u32 s1 = digests_buf[digests_offset].digest_buf[1];
-
- __local u32 s_S[64];
-
- if (lid == 0)
- {
- s_S[ 0 + vid] = -((s0 >> vid) & 1);
- }
- else if (lid == 1)
- {
- s_S[32 + vid] = -((s1 >> vid) & 1);
- }
-
- barrier (CLK_LOCAL_MEM_FENCE);
if (gid >= gid_max) return;
* main
*/
- m01500s (s_S, 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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
+ m01500s (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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
}
__kernel void m01500_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 u32 gid_max)
#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));
-static 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)
{
LUT(xAA55AA5500550055, a1, a4, a6, 0xC1)
LUT(xA55AA55AF0F5F0F5, a3, a6, xAA55AA5500550055, 0x9E)
*out4 ^= x4;
}
-static void s2 (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 s2 (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)
{
LUT(xEEEEEEEE99999999, a1, a2, a6, 0x97)
LUT(xFFFFEEEE66666666, a5, a6, xEEEEEEEE99999999, 0x67)
*out4 ^= x4;
}
-static void s3 (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 s3 (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)
{
LUT(xA50FA50FA50FA50F, a1, a3, a4, 0xC9)
LUT(xF0F00F0FF0F0F0F0, a3, a5, a6, 0x4B)
*out4 ^= x4;
}
-static void s4 (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 s4 (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)
{
LUT(x55F055F055F055F0, a1, a3, a4, 0x72)
LUT(xA500F5F0A500F5F0, a3, a5, x55F055F055F055F0, 0xAD)
*out4 ^= x4;
}
-static void s5 (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 s5 (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)
{
LUT(xA0A0A0A0FFFFFFFF, a1, a3, a6, 0xAB)
LUT(xFFFF00005555FFFF, a1, a5, a6, 0xB9)
*out4 ^= x4;
}
-static void s6 (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 s6 (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)
{
LUT(x5050F5F55050F5F5, a1, a3, a5, 0xB2)
LUT(x6363C6C66363C6C6, a1, a2, x5050F5F55050F5F5, 0x66)
*out4 ^= x4;
}
-static void s7 (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 s7 (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)
{
LUT(x88AA88AA88AA88AA, a1, a2, a4, 0x0B)
LUT(xAAAAFF00AAAAFF00, a1, a4, a5, 0x27)
*out4 ^= x4;
}
-static void s8 (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 s8 (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)
{
LUT(xEEEE3333EEEE3333, a1, a2, a5, 0x9D)
LUT(xBBBBBBBBBBBBBBBB, a1, a1, a2, 0x83)
* The effort has been sponsored by Rapid7: http://www.rapid7.com
*/
-static 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)
{
u32 x55005500, x5A0F5A0F, x3333FFFF, x66666666, x22226666, x2D2D6969,
x25202160;
*out4 ^= x31;
}
-static void s2 (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 s2 (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)
{
u32 x33CC33CC;
u32 x55550000, x00AA00FF, x33BB33FF;
*out4 ^= x31;
}
-static void s3 (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 s3 (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)
{
u32 x44444444, x0F0FF0F0, x4F4FF4F4, x00FFFF00, x00AAAA00, x4FE55EF4;
u32 x3C3CC3C3, x3C3C0000, x7373F4F4, x0C840A00;
*out3 ^= x21;
}
-static void s4 (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 s4 (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)
{
u32 x5A5A5A5A, x0F0FF0F0;
u32 x33FF33FF, x33FFCC00, x0C0030F0, x0C0CC0C0, x0CF3C03F, x5EFBDA7F,
*out4 ^= x31;
}
-static void s5 (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 s5 (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)
{
u32 x77777777, x77770000, x22225555, x11116666, x1F1F6F6F;
u32 x70700000, x43433333, x00430033, x55557777, x55167744, x5A19784B;
*out2 ^= x11;
}
-static void s6 (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 s6 (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)
{
u32 x33CC33CC;
u32 x3333FFFF, x11115555, x22DD6699, x22DD9966, x00220099;
*out1 ^= x01;
}
-static void s7 (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 s7 (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)
{
u32 x0FF00FF0, x3CC33CC3, x00003CC3, x0F000F00, x5A555A55, x00001841;
u32 x00000F00, x33333C33, x7B777E77, x0FF0F00F, x74878E78;
*out2 ^= x11;
}
-static void s8 (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 s8 (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)
{
u32 x0C0C0C0C, x0000F0F0, x00FFF00F, x00555005, x00515001;
u32 x33000330, x77555775, x30303030, x3030CFCF, x30104745, x30555745;
* The effort has been sponsored by Rapid7: http://www.rapid7.com
*/
-static 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)
{
u32 x55005500, x5A0F5A0F, x3333FFFF, x66666666, x22226666, x2D2D6969,
x25202160;
*out4 ^= x31;
}
-static void s2 (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 s2 (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)
{
u32 x33CC33CC;
u32 x55550000, x00AA00FF, x33BB33FF;
*out4 ^= x31;
}
-static void s3 (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 s3 (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)
{
u32 x44444444, x0F0FF0F0, x4F4FF4F4, x00FFFF00, x00AAAA00, x4FE55EF4;
u32 x3C3CC3C3, x3C3C0000, x7373F4F4, x0C840A00;
*out3 ^= x21;
}
-static void s4 (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 s4 (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)
{
u32 x5A5A5A5A, x0F0FF0F0;
u32 x33FF33FF, x33FFCC00, x0C0030F0, x0C0CC0C0, x0CF3C03F, x5EFBDA7F,
*out4 ^= x31;
}
-static void s5 (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 s5 (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)
{
u32 x77777777, x77770000, x22225555, x11116666, x1F1F6F6F;
u32 x70700000, x43433333, x00430033, x55557777, x55167744, x5A19784B;
*out2 ^= x11;
}
-static void s6 (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 s6 (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)
{
u32 x33CC33CC;
u32 x3333FFFF, x11115555, x22DD6699, x22DD9966, x00220099;
*out1 ^= x01;
}
-static void s7 (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 s7 (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)
{
u32 x0FF00FF0, x3CC33CC3, x00003CC3, x0F000F00, x5A555A55, x00001841;
u32 x00000F00, x33333C33, x7B777E77, x0FF0F00F, x74878E78;
*out2 ^= x11;
}
-static void s8 (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 s8 (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)
{
u32 x0C0C0C0C, x0000F0F0, x00FFF00F, x00555005, x00515001;
u32 x33000330, x77555775, x30303030, x3030CFCF, x30104745, x30555745;
#define KEYSET07 { k00 = K31; k01 = K35; k02 = K52; k03 = K43; k04 = K08; k05 = K37; k06 = K51; k07 = K15; k08 = K49; k09 = K30; k10 = K07; k11 = K02; k12 = K50; k13 = K21; k14 = K45; k15 = K44; k16 = K29; k17 = K16; k18 = K42; k19 = K23; k20 = K22; k21 = K14; k22 = K38; k23 = K01; k24 = K10; k25 = K47; k26 = K53; k27 = K11; k28 = K27; k29 = K26; k30 = K05; k31 = K17; k32 = K54; k33 = K41; k34 = K39; k35 = K20; k36 = K48; k37 = K13; k38 = K24; k39 = K19; k40 = K32; k41 = K40; k42 = K34; k43 = K03; k44 = K06; k45 = K18; k46 = K12; k47 = K46; }
#define KEYSET17 { k00 = K15; k01 = K51; k02 = K36; k03 = K02; k04 = K49; k05 = K21; k06 = K35; k07 = K31; k08 = K08; k09 = K14; k10 = K23; k11 = K43; k12 = K09; k13 = K37; k14 = K29; k15 = K28; k16 = K45; k17 = K00; k18 = K01; k19 = K07; k20 = K38; k21 = K30; k22 = K22; k23 = K42; k24 = K26; k25 = K04; k26 = K41; k27 = K54; k28 = K39; k29 = K10; k30 = K48; k31 = K33; k32 = K11; k33 = K53; k34 = K27; k35 = K32; k36 = K05; k37 = K25; k38 = K40; k39 = K03; k40 = K20; k41 = K24; k42 = K46; k43 = K19; k44 = K18; k45 = K06; k46 = K55; k47 = K34; }
-static void DES (const u32 K00, const u32 K01, const u32 K02, const u32 K03, const u32 K04, const u32 K05, const u32 K06, const u32 K07, const u32 K08, const u32 K09, const u32 K10, const u32 K11, const u32 K12, const u32 K13, const u32 K14, const u32 K15, const u32 K16, const u32 K17, const u32 K18, const u32 K19, const u32 K20, const u32 K21, const u32 K22, const u32 K23, const u32 K24, const u32 K25, const u32 K26, const u32 K27, const u32 K28, const u32 K29, const u32 K30, const u32 K31, const u32 K32, const u32 K33, const u32 K34, const u32 K35, const u32 K36, const u32 K37, const u32 K38, const u32 K39, const u32 K40, const u32 K41, const u32 K42, const u32 K43, const u32 K44, const u32 K45, const u32 K46, const u32 K47, const u32 K48, const u32 K49, const u32 K50, const u32 K51, const u32 K52, const u32 K53, const u32 K54, const u32 K55, u32 *D00, u32 *D01, u32 *D02, u32 *D03, u32 *D04, u32 *D05, u32 *D06, u32 *D07, u32 *D08, u32 *D09, u32 *D10, u32 *D11, u32 *D12, u32 *D13, u32 *D14, u32 *D15, u32 *D16, u32 *D17, u32 *D18, u32 *D19, u32 *D20, u32 *D21, u32 *D22, u32 *D23, u32 *D24, u32 *D25, u32 *D26, u32 *D27, u32 *D28, u32 *D29, u32 *D30, u32 *D31, u32 *D32, u32 *D33, u32 *D34, u32 *D35, u32 *D36, u32 *D37, u32 *D38, u32 *D39, u32 *D40, u32 *D41, u32 *D42, u32 *D43, u32 *D44, u32 *D45, u32 *D46, u32 *D47, u32 *D48, u32 *D49, u32 *D50, u32 *D51, u32 *D52, u32 *D53, u32 *D54, u32 *D55, u32 *D56, u32 *D57, u32 *D58, u32 *D59, u32 *D60, u32 *D61, u32 *D62, u32 *D63)
+void DES (const u32 K00, const u32 K01, const u32 K02, const u32 K03, const u32 K04, const u32 K05, const u32 K06, const u32 K07, const u32 K08, const u32 K09, const u32 K10, const u32 K11, const u32 K12, const u32 K13, const u32 K14, const u32 K15, const u32 K16, const u32 K17, const u32 K18, const u32 K19, const u32 K20, const u32 K21, const u32 K22, const u32 K23, const u32 K24, const u32 K25, const u32 K26, const u32 K27, const u32 K28, const u32 K29, const u32 K30, const u32 K31, const u32 K32, const u32 K33, const u32 K34, const u32 K35, const u32 K36, const u32 K37, const u32 K38, const u32 K39, const u32 K40, const u32 K41, const u32 K42, const u32 K43, const u32 K44, const u32 K45, const u32 K46, const u32 K47, const u32 K48, const u32 K49, const u32 K50, const u32 K51, const u32 K52, const u32 K53, const u32 K54, const u32 K55, u32 *D00, u32 *D01, u32 *D02, u32 *D03, u32 *D04, u32 *D05, u32 *D06, u32 *D07, u32 *D08, u32 *D09, u32 *D10, u32 *D11, u32 *D12, u32 *D13, u32 *D14, u32 *D15, u32 *D16, u32 *D17, u32 *D18, u32 *D19, u32 *D20, u32 *D21, u32 *D22, u32 *D23, u32 *D24, u32 *D25, u32 *D26, u32 *D27, u32 *D28, u32 *D29, u32 *D30, u32 *D31, u32 *D32, u32 *D33, u32 *D34, u32 *D35, u32 *D36, u32 *D37, u32 *D38, u32 *D39, u32 *D40, u32 *D41, u32 *D42, u32 *D43, u32 *D44, u32 *D45, u32 *D46, u32 *D47, u32 *D48, u32 *D49, u32 *D50, u32 *D51, u32 *D52, u32 *D53, u32 *D54, u32 *D55, u32 *D56, u32 *D57, u32 *D58, u32 *D59, u32 *D60, u32 *D61, u32 *D62, u32 *D63)
{
KXX_DECL u32 k00, k01, k02, k03, k04, k05;
KXX_DECL u32 k06, k07, k08, k09, k10, k11;
}
}
-static void transpose32c (u32 data[32])
+void transpose32c (u32 data[32])
{
#define swap(x,y,j,m) \
t = ((x) ^ ((y) >> (j))) & (m); \
swap (data[30], data[31], 1, 0x55555555);
}
-static void m03000m (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 m03000m (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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)
{
/**
* base
const u32 w0s = pws[gid].i[0];
const u32 w1s = pws[gid].i[1];
- const u32 K00 = -((w0s >> ( 0 + 7)) & 1);
- const u32 K01 = -((w0s >> ( 0 + 6)) & 1);
- const u32 K02 = -((w0s >> ( 0 + 5)) & 1);
- const u32 K03 = -((w0s >> ( 0 + 4)) & 1);
- const u32 K04 = -((w0s >> ( 0 + 3)) & 1);
- const u32 K05 = -((w0s >> ( 0 + 2)) & 1);
- const u32 K06 = -((w0s >> ( 0 + 1)) & 1);
- const u32 K07 = -((w0s >> ( 0 + 0)) & 1);
- const u32 K08 = -((w0s >> ( 8 + 7)) & 1);
- const u32 K09 = -((w0s >> ( 8 + 6)) & 1);
- const u32 K10 = -((w0s >> ( 8 + 5)) & 1);
- const u32 K11 = -((w0s >> ( 8 + 4)) & 1);
- const u32 K12 = -((w0s >> ( 8 + 3)) & 1);
- const u32 K13 = -((w0s >> ( 8 + 2)) & 1);
- const u32 K14 = -((w0s >> ( 8 + 1)) & 1);
- const u32 K15 = -((w0s >> ( 8 + 0)) & 1);
- const u32 K16 = -((w0s >> (16 + 7)) & 1);
- const u32 K17 = -((w0s >> (16 + 6)) & 1);
- const u32 K18 = -((w0s >> (16 + 5)) & 1);
- const u32 K19 = -((w0s >> (16 + 4)) & 1);
- const u32 K20 = -((w0s >> (16 + 3)) & 1);
- const u32 K21 = -((w0s >> (16 + 2)) & 1);
- const u32 K22 = -((w0s >> (16 + 1)) & 1);
- const u32 K23 = -((w0s >> (16 + 0)) & 1);
- const u32 K24 = -((w0s >> (24 + 7)) & 1);
- const u32 K25 = -((w0s >> (24 + 6)) & 1);
- const u32 K26 = -((w0s >> (24 + 5)) & 1);
- const u32 K27 = -((w0s >> (24 + 4)) & 1);
- const u32 K28 = -((w0s >> (24 + 3)) & 1);
- const u32 K29 = -((w0s >> (24 + 2)) & 1);
- const u32 K30 = -((w0s >> (24 + 1)) & 1);
- const u32 K31 = -((w0s >> (24 + 0)) & 1);
- const u32 K32 = -((w1s >> ( 0 + 7)) & 1);
- const u32 K33 = -((w1s >> ( 0 + 6)) & 1);
- const u32 K34 = -((w1s >> ( 0 + 5)) & 1);
- const u32 K35 = -((w1s >> ( 0 + 4)) & 1);
- const u32 K36 = -((w1s >> ( 0 + 3)) & 1);
- const u32 K37 = -((w1s >> ( 0 + 2)) & 1);
- const u32 K38 = -((w1s >> ( 0 + 1)) & 1);
- const u32 K39 = -((w1s >> ( 0 + 0)) & 1);
- const u32 K40 = -((w1s >> ( 8 + 7)) & 1);
- const u32 K41 = -((w1s >> ( 8 + 6)) & 1);
- const u32 K42 = -((w1s >> ( 8 + 5)) & 1);
- const u32 K43 = -((w1s >> ( 8 + 4)) & 1);
- const u32 K44 = -((w1s >> ( 8 + 3)) & 1);
- const u32 K45 = -((w1s >> ( 8 + 2)) & 1);
- const u32 K46 = -((w1s >> ( 8 + 1)) & 1);
- const u32 K47 = -((w1s >> ( 8 + 0)) & 1);
- const u32 K48 = -((w1s >> (16 + 7)) & 1);
- const u32 K49 = -((w1s >> (16 + 6)) & 1);
- const u32 K50 = -((w1s >> (16 + 5)) & 1);
- const u32 K51 = -((w1s >> (16 + 4)) & 1);
- const u32 K52 = -((w1s >> (16 + 3)) & 1);
- const u32 K53 = -((w1s >> (16 + 2)) & 1);
- const u32 K54 = -((w1s >> (16 + 1)) & 1);
- const u32 K55 = -((w1s >> (16 + 0)) & 1);
+ #define K00 -((w0s >> ( 0 + 7)) & 1)
+ #define K01 -((w0s >> ( 0 + 6)) & 1)
+ #define K02 -((w0s >> ( 0 + 5)) & 1)
+ #define K03 -((w0s >> ( 0 + 4)) & 1)
+ #define K04 -((w0s >> ( 0 + 3)) & 1)
+ #define K05 -((w0s >> ( 0 + 2)) & 1)
+ #define K06 -((w0s >> ( 0 + 1)) & 1)
+ #define K07 -((w0s >> ( 0 + 0)) & 1)
+ #define K08 -((w0s >> ( 8 + 7)) & 1)
+ #define K09 -((w0s >> ( 8 + 6)) & 1)
+ #define K10 -((w0s >> ( 8 + 5)) & 1)
+ #define K11 -((w0s >> ( 8 + 4)) & 1)
+ #define K12 -((w0s >> ( 8 + 3)) & 1)
+ #define K13 -((w0s >> ( 8 + 2)) & 1)
+ #define K14 -((w0s >> ( 8 + 1)) & 1)
+ #define K15 -((w0s >> ( 8 + 0)) & 1)
+ #define K16 -((w0s >> (16 + 7)) & 1)
+ #define K17 -((w0s >> (16 + 6)) & 1)
+ #define K18 -((w0s >> (16 + 5)) & 1)
+ #define K19 -((w0s >> (16 + 4)) & 1)
+ #define K20 -((w0s >> (16 + 3)) & 1)
+ #define K21 -((w0s >> (16 + 2)) & 1)
+ #define K22 -((w0s >> (16 + 1)) & 1)
+ #define K23 -((w0s >> (16 + 0)) & 1)
+ #define K24 -((w0s >> (24 + 7)) & 1)
+ #define K25 -((w0s >> (24 + 6)) & 1)
+ #define K26 -((w0s >> (24 + 5)) & 1)
+ #define K27 -((w0s >> (24 + 4)) & 1)
+ #define K28 -((w0s >> (24 + 3)) & 1)
+ #define K29 -((w0s >> (24 + 2)) & 1)
+ #define K30 -((w0s >> (24 + 1)) & 1)
+ #define K31 -((w0s >> (24 + 0)) & 1)
+ #define K32 -((w1s >> ( 0 + 7)) & 1)
+ #define K33 -((w1s >> ( 0 + 6)) & 1)
+ #define K34 -((w1s >> ( 0 + 5)) & 1)
+ #define K35 -((w1s >> ( 0 + 4)) & 1)
+ #define K36 -((w1s >> ( 0 + 3)) & 1)
+ #define K37 -((w1s >> ( 0 + 2)) & 1)
+ #define K38 -((w1s >> ( 0 + 1)) & 1)
+ #define K39 -((w1s >> ( 0 + 0)) & 1)
+ #define K40 -((w1s >> ( 8 + 7)) & 1)
+ #define K41 -((w1s >> ( 8 + 6)) & 1)
+ #define K42 -((w1s >> ( 8 + 5)) & 1)
+ #define K43 -((w1s >> ( 8 + 4)) & 1)
+ #define K44 -((w1s >> ( 8 + 3)) & 1)
+ #define K45 -((w1s >> ( 8 + 2)) & 1)
+ #define K46 -((w1s >> ( 8 + 1)) & 1)
+ #define K47 -((w1s >> ( 8 + 0)) & 1)
+ #define K48 -((w1s >> (16 + 7)) & 1)
+ #define K49 -((w1s >> (16 + 6)) & 1)
+ #define K50 -((w1s >> (16 + 5)) & 1)
+ #define K51 -((w1s >> (16 + 4)) & 1)
+ #define K52 -((w1s >> (16 + 3)) & 1)
+ #define K53 -((w1s >> (16 + 2)) & 1)
+ #define K54 -((w1s >> (16 + 1)) & 1)
+ #define K55 -((w1s >> (16 + 0)) & 1)
/**
* inner loop
}
}
-static void m03000s (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 m03000s (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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)
{
/**
* base
* digest
*/
- const u32 s0 = digests_buf[digests_offset].digest_buf[0];
- const u32 s1 = digests_buf[digests_offset].digest_buf[1];
-
- const u32 S00 = -((s0 >> 0) & 1);
- const u32 S01 = -((s0 >> 1) & 1);
- const u32 S02 = -((s0 >> 2) & 1);
- const u32 S03 = -((s0 >> 3) & 1);
- const u32 S04 = -((s0 >> 4) & 1);
- const u32 S05 = -((s0 >> 5) & 1);
- const u32 S06 = -((s0 >> 6) & 1);
- const u32 S07 = -((s0 >> 7) & 1);
- const u32 S08 = -((s0 >> 8) & 1);
- const u32 S09 = -((s0 >> 9) & 1);
- const u32 S10 = -((s0 >> 10) & 1);
- const u32 S11 = -((s0 >> 11) & 1);
- const u32 S12 = -((s0 >> 12) & 1);
- const u32 S13 = -((s0 >> 13) & 1);
- const u32 S14 = -((s0 >> 14) & 1);
- const u32 S15 = -((s0 >> 15) & 1);
- const u32 S16 = -((s0 >> 16) & 1);
- const u32 S17 = -((s0 >> 17) & 1);
- const u32 S18 = -((s0 >> 18) & 1);
- const u32 S19 = -((s0 >> 19) & 1);
- const u32 S20 = -((s0 >> 20) & 1);
- const u32 S21 = -((s0 >> 21) & 1);
- const u32 S22 = -((s0 >> 22) & 1);
- const u32 S23 = -((s0 >> 23) & 1);
- const u32 S24 = -((s0 >> 24) & 1);
- const u32 S25 = -((s0 >> 25) & 1);
- const u32 S26 = -((s0 >> 26) & 1);
- const u32 S27 = -((s0 >> 27) & 1);
- const u32 S28 = -((s0 >> 28) & 1);
- const u32 S29 = -((s0 >> 29) & 1);
- const u32 S30 = -((s0 >> 30) & 1);
- const u32 S31 = -((s0 >> 31) & 1);
- const u32 S32 = -((s1 >> 0) & 1);
- const u32 S33 = -((s1 >> 1) & 1);
- const u32 S34 = -((s1 >> 2) & 1);
- const u32 S35 = -((s1 >> 3) & 1);
- const u32 S36 = -((s1 >> 4) & 1);
- const u32 S37 = -((s1 >> 5) & 1);
- const u32 S38 = -((s1 >> 6) & 1);
- const u32 S39 = -((s1 >> 7) & 1);
- const u32 S40 = -((s1 >> 8) & 1);
- const u32 S41 = -((s1 >> 9) & 1);
- const u32 S42 = -((s1 >> 10) & 1);
- const u32 S43 = -((s1 >> 11) & 1);
- const u32 S44 = -((s1 >> 12) & 1);
- const u32 S45 = -((s1 >> 13) & 1);
- const u32 S46 = -((s1 >> 14) & 1);
- const u32 S47 = -((s1 >> 15) & 1);
- const u32 S48 = -((s1 >> 16) & 1);
- const u32 S49 = -((s1 >> 17) & 1);
- const u32 S50 = -((s1 >> 18) & 1);
- const u32 S51 = -((s1 >> 19) & 1);
- const u32 S52 = -((s1 >> 20) & 1);
- const u32 S53 = -((s1 >> 21) & 1);
- const u32 S54 = -((s1 >> 22) & 1);
- const u32 S55 = -((s1 >> 23) & 1);
- const u32 S56 = -((s1 >> 24) & 1);
- const u32 S57 = -((s1 >> 25) & 1);
- const u32 S58 = -((s1 >> 26) & 1);
- const u32 S59 = -((s1 >> 27) & 1);
- const u32 S60 = -((s1 >> 28) & 1);
- const u32 S61 = -((s1 >> 29) & 1);
- const u32 S62 = -((s1 >> 30) & 1);
- const u32 S63 = -((s1 >> 31) & 1);
+ const u32 s0 = digests_buf[0].digest_buf[0];
+ const u32 s1 = digests_buf[0].digest_buf[1];
+
+ #define S00 -((s0 >> 0) & 1)
+ #define S01 -((s0 >> 1) & 1)
+ #define S02 -((s0 >> 2) & 1)
+ #define S03 -((s0 >> 3) & 1)
+ #define S04 -((s0 >> 4) & 1)
+ #define S05 -((s0 >> 5) & 1)
+ #define S06 -((s0 >> 6) & 1)
+ #define S07 -((s0 >> 7) & 1)
+ #define S08 -((s0 >> 8) & 1)
+ #define S09 -((s0 >> 9) & 1)
+ #define S10 -((s0 >> 10) & 1)
+ #define S11 -((s0 >> 11) & 1)
+ #define S12 -((s0 >> 12) & 1)
+ #define S13 -((s0 >> 13) & 1)
+ #define S14 -((s0 >> 14) & 1)
+ #define S15 -((s0 >> 15) & 1)
+ #define S16 -((s0 >> 16) & 1)
+ #define S17 -((s0 >> 17) & 1)
+ #define S18 -((s0 >> 18) & 1)
+ #define S19 -((s0 >> 19) & 1)
+ #define S20 -((s0 >> 20) & 1)
+ #define S21 -((s0 >> 21) & 1)
+ #define S22 -((s0 >> 22) & 1)
+ #define S23 -((s0 >> 23) & 1)
+ #define S24 -((s0 >> 24) & 1)
+ #define S25 -((s0 >> 25) & 1)
+ #define S26 -((s0 >> 26) & 1)
+ #define S27 -((s0 >> 27) & 1)
+ #define S28 -((s0 >> 28) & 1)
+ #define S29 -((s0 >> 29) & 1)
+ #define S30 -((s0 >> 30) & 1)
+ #define S31 -((s0 >> 31) & 1)
+ #define S32 -((s1 >> 0) & 1)
+ #define S33 -((s1 >> 1) & 1)
+ #define S34 -((s1 >> 2) & 1)
+ #define S35 -((s1 >> 3) & 1)
+ #define S36 -((s1 >> 4) & 1)
+ #define S37 -((s1 >> 5) & 1)
+ #define S38 -((s1 >> 6) & 1)
+ #define S39 -((s1 >> 7) & 1)
+ #define S40 -((s1 >> 8) & 1)
+ #define S41 -((s1 >> 9) & 1)
+ #define S42 -((s1 >> 10) & 1)
+ #define S43 -((s1 >> 11) & 1)
+ #define S44 -((s1 >> 12) & 1)
+ #define S45 -((s1 >> 13) & 1)
+ #define S46 -((s1 >> 14) & 1)
+ #define S47 -((s1 >> 15) & 1)
+ #define S48 -((s1 >> 16) & 1)
+ #define S49 -((s1 >> 17) & 1)
+ #define S50 -((s1 >> 18) & 1)
+ #define S51 -((s1 >> 19) & 1)
+ #define S52 -((s1 >> 20) & 1)
+ #define S53 -((s1 >> 21) & 1)
+ #define S54 -((s1 >> 22) & 1)
+ #define S55 -((s1 >> 23) & 1)
+ #define S56 -((s1 >> 24) & 1)
+ #define S57 -((s1 >> 25) & 1)
+ #define S58 -((s1 >> 26) & 1)
+ #define S59 -((s1 >> 27) & 1)
+ #define S60 -((s1 >> 28) & 1)
+ #define S61 -((s1 >> 29) & 1)
+ #define S62 -((s1 >> 30) & 1)
+ #define S63 -((s1 >> 31) & 1)
/**
* base
const u32 w0s = pws[gid].i[0];
const u32 w1s = pws[gid].i[1];
- const u32 K00 = -((w0s >> ( 0 + 7)) & 1);
- const u32 K01 = -((w0s >> ( 0 + 6)) & 1);
- const u32 K02 = -((w0s >> ( 0 + 5)) & 1);
- const u32 K03 = -((w0s >> ( 0 + 4)) & 1);
- const u32 K04 = -((w0s >> ( 0 + 3)) & 1);
- const u32 K05 = -((w0s >> ( 0 + 2)) & 1);
- const u32 K06 = -((w0s >> ( 0 + 1)) & 1);
- const u32 K07 = -((w0s >> ( 0 + 0)) & 1);
- const u32 K08 = -((w0s >> ( 8 + 7)) & 1);
- const u32 K09 = -((w0s >> ( 8 + 6)) & 1);
- const u32 K10 = -((w0s >> ( 8 + 5)) & 1);
- const u32 K11 = -((w0s >> ( 8 + 4)) & 1);
- const u32 K12 = -((w0s >> ( 8 + 3)) & 1);
- const u32 K13 = -((w0s >> ( 8 + 2)) & 1);
- const u32 K14 = -((w0s >> ( 8 + 1)) & 1);
- const u32 K15 = -((w0s >> ( 8 + 0)) & 1);
- const u32 K16 = -((w0s >> (16 + 7)) & 1);
- const u32 K17 = -((w0s >> (16 + 6)) & 1);
- const u32 K18 = -((w0s >> (16 + 5)) & 1);
- const u32 K19 = -((w0s >> (16 + 4)) & 1);
- const u32 K20 = -((w0s >> (16 + 3)) & 1);
- const u32 K21 = -((w0s >> (16 + 2)) & 1);
- const u32 K22 = -((w0s >> (16 + 1)) & 1);
- const u32 K23 = -((w0s >> (16 + 0)) & 1);
- const u32 K24 = -((w0s >> (24 + 7)) & 1);
- const u32 K25 = -((w0s >> (24 + 6)) & 1);
- const u32 K26 = -((w0s >> (24 + 5)) & 1);
- const u32 K27 = -((w0s >> (24 + 4)) & 1);
- const u32 K28 = -((w0s >> (24 + 3)) & 1);
- const u32 K29 = -((w0s >> (24 + 2)) & 1);
- const u32 K30 = -((w0s >> (24 + 1)) & 1);
- const u32 K31 = -((w0s >> (24 + 0)) & 1);
- const u32 K32 = -((w1s >> ( 0 + 7)) & 1);
- const u32 K33 = -((w1s >> ( 0 + 6)) & 1);
- const u32 K34 = -((w1s >> ( 0 + 5)) & 1);
- const u32 K35 = -((w1s >> ( 0 + 4)) & 1);
- const u32 K36 = -((w1s >> ( 0 + 3)) & 1);
- const u32 K37 = -((w1s >> ( 0 + 2)) & 1);
- const u32 K38 = -((w1s >> ( 0 + 1)) & 1);
- const u32 K39 = -((w1s >> ( 0 + 0)) & 1);
- const u32 K40 = -((w1s >> ( 8 + 7)) & 1);
- const u32 K41 = -((w1s >> ( 8 + 6)) & 1);
- const u32 K42 = -((w1s >> ( 8 + 5)) & 1);
- const u32 K43 = -((w1s >> ( 8 + 4)) & 1);
- const u32 K44 = -((w1s >> ( 8 + 3)) & 1);
- const u32 K45 = -((w1s >> ( 8 + 2)) & 1);
- const u32 K46 = -((w1s >> ( 8 + 1)) & 1);
- const u32 K47 = -((w1s >> ( 8 + 0)) & 1);
- const u32 K48 = -((w1s >> (16 + 7)) & 1);
- const u32 K49 = -((w1s >> (16 + 6)) & 1);
- const u32 K50 = -((w1s >> (16 + 5)) & 1);
- const u32 K51 = -((w1s >> (16 + 4)) & 1);
- const u32 K52 = -((w1s >> (16 + 3)) & 1);
- const u32 K53 = -((w1s >> (16 + 2)) & 1);
- const u32 K54 = -((w1s >> (16 + 1)) & 1);
- const u32 K55 = -((w1s >> (16 + 0)) & 1);
+ #define K00 -((w0s >> ( 0 + 7)) & 1)
+ #define K01 -((w0s >> ( 0 + 6)) & 1)
+ #define K02 -((w0s >> ( 0 + 5)) & 1)
+ #define K03 -((w0s >> ( 0 + 4)) & 1)
+ #define K04 -((w0s >> ( 0 + 3)) & 1)
+ #define K05 -((w0s >> ( 0 + 2)) & 1)
+ #define K06 -((w0s >> ( 0 + 1)) & 1)
+ #define K07 -((w0s >> ( 0 + 0)) & 1)
+ #define K08 -((w0s >> ( 8 + 7)) & 1)
+ #define K09 -((w0s >> ( 8 + 6)) & 1)
+ #define K10 -((w0s >> ( 8 + 5)) & 1)
+ #define K11 -((w0s >> ( 8 + 4)) & 1)
+ #define K12 -((w0s >> ( 8 + 3)) & 1)
+ #define K13 -((w0s >> ( 8 + 2)) & 1)
+ #define K14 -((w0s >> ( 8 + 1)) & 1)
+ #define K15 -((w0s >> ( 8 + 0)) & 1)
+ #define K16 -((w0s >> (16 + 7)) & 1)
+ #define K17 -((w0s >> (16 + 6)) & 1)
+ #define K18 -((w0s >> (16 + 5)) & 1)
+ #define K19 -((w0s >> (16 + 4)) & 1)
+ #define K20 -((w0s >> (16 + 3)) & 1)
+ #define K21 -((w0s >> (16 + 2)) & 1)
+ #define K22 -((w0s >> (16 + 1)) & 1)
+ #define K23 -((w0s >> (16 + 0)) & 1)
+ #define K24 -((w0s >> (24 + 7)) & 1)
+ #define K25 -((w0s >> (24 + 6)) & 1)
+ #define K26 -((w0s >> (24 + 5)) & 1)
+ #define K27 -((w0s >> (24 + 4)) & 1)
+ #define K28 -((w0s >> (24 + 3)) & 1)
+ #define K29 -((w0s >> (24 + 2)) & 1)
+ #define K30 -((w0s >> (24 + 1)) & 1)
+ #define K31 -((w0s >> (24 + 0)) & 1)
+ #define K32 -((w1s >> ( 0 + 7)) & 1)
+ #define K33 -((w1s >> ( 0 + 6)) & 1)
+ #define K34 -((w1s >> ( 0 + 5)) & 1)
+ #define K35 -((w1s >> ( 0 + 4)) & 1)
+ #define K36 -((w1s >> ( 0 + 3)) & 1)
+ #define K37 -((w1s >> ( 0 + 2)) & 1)
+ #define K38 -((w1s >> ( 0 + 1)) & 1)
+ #define K39 -((w1s >> ( 0 + 0)) & 1)
+ #define K40 -((w1s >> ( 8 + 7)) & 1)
+ #define K41 -((w1s >> ( 8 + 6)) & 1)
+ #define K42 -((w1s >> ( 8 + 5)) & 1)
+ #define K43 -((w1s >> ( 8 + 4)) & 1)
+ #define K44 -((w1s >> ( 8 + 3)) & 1)
+ #define K45 -((w1s >> ( 8 + 2)) & 1)
+ #define K46 -((w1s >> ( 8 + 1)) & 1)
+ #define K47 -((w1s >> ( 8 + 0)) & 1)
+ #define K48 -((w1s >> (16 + 7)) & 1)
+ #define K49 -((w1s >> (16 + 6)) & 1)
+ #define K50 -((w1s >> (16 + 5)) & 1)
+ #define K51 -((w1s >> (16 + 4)) & 1)
+ #define K52 -((w1s >> (16 + 3)) & 1)
+ #define K53 -((w1s >> (16 + 2)) & 1)
+ #define K54 -((w1s >> (16 + 1)) & 1)
+ #define K55 -((w1s >> (16 + 0)) & 1)
/**
* inner loop
const u32 w0 = mod[gid];
-
for (int i = 0; i < 32; i += 8)
{
atomic_or (&words_buf_r[block].b[i + 0], (((w0 >> (i + 7)) & 1) << slice));
}
}
-__kernel void m03000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 u32 gid_max)
+__kernel void m03000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 u32 gid_max)
{
/**
* base
m03000m (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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
}
-__kernel void m03000_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 u32 gid_max)
+__kernel void m03000_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 u32 gid_max)
{
}
-__kernel void m03000_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 u32 gid_max)
+__kernel void m03000_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 u32 gid_max)
{
}
-__kernel void m03000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 u32 gid_max)
+__kernel void m03000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 u32 gid_max)
{
/**
* base
m03000s (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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
}
-__kernel void m03000_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 u32 gid_max)
+__kernel void m03000_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 u32 gid_max)
{
}
-__kernel void m03000_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 u32 gid_max)
+__kernel void m03000_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 u32 gid_max)
{
}
0x8000000000008080, 0x0000000080000001, 0x8000000080008008
};
-__constant u8 keccakf_rotc[24] =
-{
- 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
- 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
-};
-
-__constant u8 keccakf_piln[24] =
-{
- 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
- 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
-};
-
#ifndef KECCAK_ROUNDS
#define KECCAK_ROUNDS 24
#endif
const u32 pw_len = pws[gid].pw_len;
+ /**
+ * constants
+ */
+
+ const u8 keccakf_rotc[24] =
+ {
+ 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
+ 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
+ };
+
+ const u8 keccakf_piln[24] =
+ {
+ 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
+ 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
+ };
+
/**
* 0x80 keccak, very special
*/
const u32 pw_len = pws[gid].pw_len;
+ /**
+ * constants
+ */
+
+ const u8 keccakf_rotc[24] =
+ {
+ 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
+ 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
+ };
+
+ const u8 keccakf_piln[24] =
+ {
+ 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
+ 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
+ };
+
/**
* 0x80 keccak, very special
*/
0x8000000000008080, 0x0000000080000001, 0x8000000080008008
};
-__constant u8 keccakf_rotc[24] =
-{
- 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
- 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
-};
-
-__constant u8 keccakf_piln[24] =
-{
- 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
- 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
-};
-
#ifndef KECCAK_ROUNDS
#define KECCAK_ROUNDS 24
#endif
const u32 pw_l_len = pws[gid].pw_len;
+ /**
+ * constants
+ */
+
+ const u8 keccakf_rotc[24] =
+ {
+ 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
+ 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
+ };
+
+ const u8 keccakf_piln[24] =
+ {
+ 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
+ 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
+ };
+
/**
* 0x80 keccak, very special
*/
const u32 pw_l_len = pws[gid].pw_len;
+ /**
+ * constants
+ */
+
+ const u8 keccakf_rotc[24] =
+ {
+ 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
+ 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
+ };
+
+ const u8 keccakf_piln[24] =
+ {
+ 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
+ 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
+ };
+
/**
* 0x80 keccak, very special
*/
0x8000000000008080, 0x0000000080000001, 0x8000000080008008
};
-__constant u8 keccakf_rotc[24] =
-{
- 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
- 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
-};
-
-__constant u8 keccakf_piln[24] =
-{
- 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
- 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
-};
-
#ifndef KECCAK_ROUNDS
#define KECCAK_ROUNDS 24
#endif
const u32 gid = get_global_id (0);
const u32 lid = get_local_id (0);
+ /**
+ * constants
+ */
+
+ const u8 keccakf_rotc[24] =
+ {
+ 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
+ 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
+ };
+
+ const u8 keccakf_piln[24] =
+ {
+ 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
+ 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
+ };
+
/**
* 0x80 keccak, very special
*/
const u32 gid = get_global_id (0);
const u32 lid = get_local_id (0);
+ /**
+ * constants
+ */
+
+ const u8 keccakf_rotc[24] =
+ {
+ 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
+ 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
+ };
+
+ const u8 keccakf_piln[24] =
+ {
+ 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
+ 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
+ };
+
/**
* 0x80 keccak, very special
*/
Desc.: Use a fixed path for the potfile rather than session depending
Issue: 99
+Type.: Change
+File.: Rules
+Desc.: Removed some duplicate rules in T0XlCv1, d3ad0ne and dive. All rules were found due to superfluous white space differences
+Issue: 301
+
Type.: Bug
File.: Host
Desc.: Fixed a bug in combination of --restore and a user immediately aborting the session after restart
uint exec_pos;
double exec_ms[EXEC_CACHE];
+ // this is "average" speed, we'll use this for benchmark and final status screen
+
+ u64 speed_cnt_total;
+ double speed_ms_total;
+
+ // this is "current" speed
+
uint speed_pos;
u64 speed_cnt[SPEED_CACHE];
- float speed_ms[SPEED_CACHE];
+ double speed_ms[SPEED_CACHE];
+
+ // speed_rec is what additionally limits the "current" speed in time, not array elements
hc_timer_t speed_rec[SPEED_CACHE];
hc_timer_t timer_running; // timer on current dict
hc_timer_t timer_paused; // timer on current dict
- float ms_paused; // timer on current dict
+ double ms_paused; // timer on current dict
/**
* hash_info and username
i10i31i52
i10i31i52i73
i10i31i52i73i94
-i11i31
-i11i31i51
-i11i31i51i71
-i11i31i51i71i91
ci11i32
ci11i32i53
ci10i31i52
ci10i31i52i73
ci10i31i52i73i94
-ci11i31
-ci11i31i51
-ci11i31i51i71
-ci11i31i51i71i91
ui11i32
ui11i32i53
ui10i31i52
ui10i31i52i73
ui10i31i52i73i94
-ui11i31
-ui11i31i51
-ui11i31i51i71
-ui11i31i51i71i91
$1$2$a$b
$1$2$3$a$b$c
u^#
u$&
u^&
-u$*
-u^*
u$?
u^?
u$_
i1.i3.
cT1T2i1.i3.
-cT1i1.
-cT2i2.
-cT3i3.
-cT4i4.
-cT5i5.
$0d]
$1d]
so0se3si1
sa@se3si1
sa@so0si!
-sa@so0se3
so0se3si!
sa@se3si!
sa@so0ss$
cso0se3si1
csa@se3si1
csa@so0si!
-csa@so0se3
cso0se3si!
csa@se3si!
csa@so0ss$
}}}}}'5
}}}}'4
}}}}
-}}}:
}}}
}}D1{
}}
x23p1
x23d
x23
-x21:
x21
x14p2
x14dp2'6
x03
x02{
x02d
-x02+2:
+x02+2
x02
u$~$!$@$#$$
u$z
tt
t
ss5sa@se3si1so0
-ss5sa@sa4se3sl7si1si|so0
+ss5sa@se3sl7si1so0
ss5sa4sl7si|so0
ss5
so0se3
si|
si1se3
si1
-si1
se3
sck
sae
l$!$!$!
l$!$!
l$!
-l $s
-l $e $d
-l $1$2$3
l
kz5x06
krKz3x13
i61 i73
i61 i72 i85
i61 i72 i84
-i61 i72 i83
i61 i72 i82
i61 i72 i81
i61 i72 i80
i0
f'4y4
dx27f'4y4
-d:
d$~$!$@$#$$
d$z
d$y
c$2$!
c$2$
c$2
-c$1:
+c$1
c$1$~
c$1$}
c$1$|
c$1$"
c$1$!
c$1$
-c$1
c$0$~
c$0$}
c$0$|
c$ $1
c$ $0
c$
-c $1
-c $0 $1
-c $0
c
^z
^y^m
^e^n^o^t^s
^e^n^i^l
^e^l^e^t
-^e^h^t:
^e^h^t
^e^e^r^f
^e^e^n^k
]$a
]$1
] { [ }
-] ] ] ]
-] ] ] $y
] ] ] $p $h $p $b $b
-] ] ] $o
-] ] ] $m $a $n
] ] ] $d $o $g
-] ] ]
] ] K
-] ] $s
-] ] $i $e
-] ] $e $r
-] ] $a
-] ]
] $0
]
[^z
[^i
[^h
[^g
-[^f:
[^f
[^d
[^c
[[[[
[[[$1$2$3
[[[$1
-[[[ $1
-[[[ $1
-[[[
-[[
-[ ^k
-[ ]
[ [ [
[ [
[ $1
T0]
T0
Kx32p2*A5'4
-K:
K*54x31'4
K
E$sT0
E
-D5:
D5
D4
D3D2
-D3:
D3
D2]
-D2D4:
+D2D4
D2D3
D2D2z3x13
D2D1
D1D3D5
D1
D0D3
-D0 D3
C$~$!$@$#$$
C$z
C$y
C$a$b$l$e
C$a
C$S
-:x03l
-:so0
+x03l
:
.3$J'4y4
.2
'7$!$!
'7$!
'7
-'7
'6$~
'6$}
'6$|
'6$!$!
'6$!
'6
-'6
'5[
'5D3
'5D2p1
'5D2d
'5$1
'5$.
-'5 $1
-'5
'5
'4y4
'4p1
'4d
'4$y
'4$.
-'4 p1
-'4 d
-'4
'4
'3d
'3$1$2$3
'3$.
-'3 d
-'3
'3
'2$.
$~$!$@$#$$
$e$e
$e$d
$e$a$u
-$e $d
$e
$d$o$m$e
$d$o$g
$9$2
$9$1
$9$0
-$9 $9 $9
-$9 $9
$9
$8z3x13
$8$9
$8$2
$8$1
$8$0
-$8 $8
$8
$7Y1Kz3x13
$7$9
$7$2
$7$1
$7$0
-$7 $7
$7
$6p4'5D1
$6]]]
$6$2
$6$1
$6$0
-$6 $9
-$6 $6 $6
$6
$5$9
$5$8
$2$0$1$0$#
$2$0$1$0$!
$2$0
-$2 $3
-$2 $2
-$2 $1
$2 $0 $0 $2
$2 $0 $0 $0
$2
$1$5
$1$4
$1$3
-$1$2:
+$1$2
$1$2$3z3x13
$1$2$3$4$5$6$7$8$9$0
$1$2$3$4$5$6$7$8$9
$1$2$3$4$5
$1$2$3$4
$1$2$3
-$1$2
$1$1$1$1$1$1$1$1$1$1$1$1
$1$1$1$1$1$1$1$1$1$1$1
$1$1$1$1$1$1$1$1$1$1
$1$1$1$1
$1$1$1
$1$1
-$1$0:
$1$0
-$1 $3
-$1 $2 $3
-$1 $2
-$1 $1
$1
$0Z1
-$0Y1:
+$0Y1
$0*44
$0$9
$0$8
$0$1
$0$0$7}[
$0$0
-$0 $1
-$0 $0
$0
$.$z$w
$.$z$m
^\$\
^$$$
^%$%
-^&$&
^-$-
^_$_
^=$=
-^+$+
\ No newline at end of file
+^+$+
$/{
$/}z3
$0
-$0
-$0 $1
+$0$1
$0$0
$0$0$0
$0$0$1
$0$0$7
-$0$1
$0$2
$0$3
$0$4
$0z1
$0z2
$1
-$1
-$1 $2
-$1 $2 $3
+$1$2
+$1$2$3
$1$!
$1$0
$1$1
$1$1$1
-$1$2
-$1$2$3
$1$2$3$!$@$#
$1$2$3$4
$1$2$3$4$5
$1z2
$1z3
$2
-$2
-$2 $0 $0 $0
-$2 $0 $0 $1
-$2 $0 $0 $2
-$2 $0 $0 $3
-$2 $0 $0 $4
-$2 $0 $0 $5
-$2 $0 $0 $6
-$2 $0 $0 $7
-$2 $0 $0 $8
-$2 $0 $0 $9
-$2 $0 $1 $0
-$2 $0 $1 $1
-$2 $0 $1 $2
$2$!
$2$"
$2$%
$2z4
$2{
$3
-$3
$3$%
$3$*
$3$0
$3z2
$3{
$4
-$4
$4$0
$4$0$0$0
$4$1
$6z4
$6{
$7
-$7
$7$0
$7$0$0$0
$7$1
$8z2
$8z3
$8{
-$9
$9
$9$%
$9$0
ss5 se3
[ [ [
KT0
-R0se5s5&
+R0se&
i54
} $e
T0 T7 T1
T0 T7 T3
c se3 so0
+5+6
-c so0 se3
T0 T2 T7
T0 T7 T2
R5+B
^f
s1,x39$i
R9
-si1 se3
+7D1
*78D4
i70 i70
*7A}
c $1 $9 $8 $7
} o6u
-so0 se3
+5 +5 +5
[ ^v
.Ax37T7
$8xA5
^1 ^1
sa@ se3 si1
-sa@ si1 se3
*56^jT6
+8*8Bs*r
{ o41
*05^j
$4s3/
ss5 si1 se3
-ss5 se3 si1
sLax61
$#
o5n
+9swS
$d $o
ss5 se3 sa@
-ss5 sa@ se3
o6-
*51o5u
o0c
ss5 se3 so0
$1xAB
*10,1
-ss5 so0 se3
$9 $6 $6
c '8 $5
*84
sa@ si1 so0
*31 o3p *35
xB6
-sa@ so0 si1
$e $2
o7-
.Bu
^7 ^4
+5-8
u '9 $4
-sa@ so0 se3
s#=^B
c '9 $0
+8 ,9
c $1 $9 $7 $0
s`qi0L+9
se3 so0 si1
-se3 si1 so0
$y{
'AiBkc
$e $c
D8 x73
ss5 sa@ se3 si1
y1}
-ss5 si1 sa@ se3
l $p $a $s $s
-ss5 se3 sa@ si1
-ss5 se3 si1 sa@
-ss5 si1 se3 sa@
-ss5 sa@ si1 se3
^8 ^4
o0j
$4 $1 $5
o1d
ss5 sa@ si1
o69
-ss5 si1 sa@
i71 i89 i99 iA4
$b $a $y
i69
z1k
'8 $1
u $1 $9 $7 $2
-ss5 si1 so0
c '6 $9
$t $o $o
$1 DB
o2@
ss5 si1 so0 se3
saUx6BD7
-ss5 se3 si1 so0
-ss5 so0 se3 si1
szis6/*79
$f $l $y
-ss5 se3 so0 si1
-ss5 si1 se3 so0
-ss5 so0 si1 se3
snCZ1
^2 ^6
^2 ^1 ^1
Z1xB9
sa@ se3 so0 si1
^3 ^2 ^2
-sa@ si1 so0 se3
x43o5q
^2 ^3 ^1
x47T3
oAi
s0]i7b
l $9 $8
-sa@ so0 se3 si1
-sa@ se3 si1 so0
c $6 $1 $9
+7 -8 -8
-sa@ so0 si1 se3
-sa@ si1 se3 so0
^hsQb
$(
x69*A4$$
$1 $5 $7
l $5 $9
c$1$2$3$*
-s\3s\F[
+s\3[
$b $i $t
K.2
^nx82x75
*6A^z
'Ay3
ss5 sa@ so0
-ss5 so0 sa@
T3*B5
^0 ^3 ^3
c $4 $0 $0
r $1 $4
ss5 se3 so0 sa@
*6A*74
-ss5 sa@ se3 so0
o1E
x83*B0d
-ss5 sa@ so0 se3
-ss5 so0 se3 sa@
-ss5 se3 sa@ so0
$U
i5i
xA8i66xB6
$1 $8 $8
-ss5 so0 sa@ se3
p1x52
^7 ^1 ^0 ^2
C {
u 'A
ss5 si1 so0 sa@
$A i0a
-ss5 so0 sa@ si1
i4@
'B,9
*24.BY1
$2 xA2
-ss5 sa@ so0 si1
$m $e $t
+5*5B
+7 -3
-ss5 so0 si1 sa@
i1b i1e
-ss5 si1 sa@ so0
^2 ^5 ^0
-ss5 sa@ si1 so0
x4BT1]
l $6
y4u
^E i1l
c ^8 ^0
^l x17
-ss5 sa@ so0 si1 se3
-ss5 se3 so0 sa@ si1
-ss5 so0 sa@ se3 si1
s45
'8 $2
-ss5 se3 si1 sa@ so0
-ss5 so0 se3 si1 sa@
-ss5 si1 so0 sa@ se3
-ss5 sa@ so0 se3 si1
-ss5 sa@ si1 so0 se3
Y5 x47 +2
-ss5 si1 sa@ so0 se3
*80T2*5B
-ss5 sa@ se3 so0 si1
*68oA'
-ss5 si1 se3 sa@ so0
i4a x31
-ss5 si1 so0 se3 sa@
-ss5 se3 so0 si1 sa@
c '6 $0
[ K
-ss5 se3 sa@ si1 so0
-ss5 so0 sa@ si1 se3
o2r [
-ss5 se3 si1 so0 sa@
-ss5 so0 si1 se3 sa@
+B -D
Y5 +2 x47
$3 $5 $7
*50 y3
-ss5 si1 se3 so0 sa@
]-8
-ss5 so0 si1 sa@ se3
-ss5 si1 sa@ se3 so0
-ss5 sa@ si1 se3 so0
-ss5 se3 sa@ so0 si1
-ss5 so0 se3 sa@ si1
^ysVn
s'{x8A*32
^f }
ls10
$l Z1
^X $X $x ^x
-s.hx13sh<
+s.<x13
^X ^x $X $x
i6v
s1!
fiAD'7
+5 swV
$4 +9
-c se3 si1
x87i78iB)
$-iB,x96
$5 $0 $3
i2M
p4 o9e
x61 $8
-slOsOZ
+slZ
$! ^1
^0 ^1 ^8
i62 i70 i81 i92
d '8 $5
x36+0sT6
x03 q
-c so0 si1
$1o3c
o6N
x27x95$A
$4 $0 $9
^9 ^4 ^3
i56 i58
-si1 sa@
$2*7AT6
o4c x02
$h $u $m
-4sC!
*02 o1a
s-So92$S
-sa4 se3 si1
i5o i6u
$5 $7 $6
$8 $0 $6
*13 *04
x23^K
$2 x61
-sk@s@{oB8
+sk{oB8
^8 ^1 ^3
o0r c
[ o2z
$8 $7 $5
$8 $7 $4
^f ^4
-se3 sa@
x0Ao3fr
$D'B
+0 $s
^1 ^3 ^8
i64 i72 i80
s{D*58r
-+8sPMsMW
++8sPW
$t $h
x37sFp
^8 ^3 ^9
i52 c
o74 *97
^A c
-c sa4 so0
scfT6xB9
*02 x42
oBli6mx91
x17sbl
T5 T6 T7 T8
x21 ^r
-so0 si1
s0I,6
i2s k
i7c i7e
-0 ,3
i0n i1o
i1o i2l i3l
-sa4 so0 se3
o4c +4
i9!
,0i6h
,1s}%
qx18sQX
+0 -9 +1
-sa4 si1 so0
xABx34
x91i2k
$1 $1 $0 $9
qs{cx48
+1 x23
*B7^g
-so0 se3 si1
$9 $9 $! $!
$k $e $g
[ $.
+2K
c $x ^x $X ^X
o4(
-so0 si1 se3
Y5KxB7
-Ao6h
+A +A -8
u d $3
+7 i7_
]{}
-so0 sa4
i5i i6n
iBd
*80i72
i01 i19 i28 i37
+3 +3 +4
sKMx0A
-sYfsfut
+sYut
$1 ^1
-0 o1j
x71 T1
i2y l
x56s `R8
i6f *64
-c so0 sa@
$x Z1
k +5
x43q
*1B$5Z2
[*09Z2
Z4 x38
-so0 sa@
LB
*47*80
x84 x98
D0 r
+1 l
o48 K
-sTDo5OsT*
+sTDo5O
t ,3
o1b o0r
x54 p3
*B0 x07
-4 +2 $A
x02 +6
-so! $= s!=
+so=$=
i74 i5a
o0k *46
i7* x62 +3
i13 i34
-so! s!= $=
$v $i $m
^. *10
x08 z1
suiD8
^n D1 x61
$ksy#
-c se3 sa@
sdy
DA $v
x67Y3D1
^1 z2
i01 i19 i28 i34
s2Y
-c so0 sl1
i65 '9
o5Y k
i3E K
s*hxA1z3
iB*
y1 x03
+$ {
sbR
oAv
i15 i35
K -0
i76 i7e
x71Y4
+^}
-4 i5d
-7^k
D3 i24
i70 i7d
soEo68[
xA2 p2 K
-c sl1 se3
x0B,B+3
xA9+8
*ABo2R
*B7i7L
tZ1
o1;
-c sa@ si1
D7 o3d
*B1i0r
i4$
*10 i1b
T9sApsu2
+7 o46
-c se3 so0 sa@
i7+ x53 $O
$Qt
x48d'8
sOZ^9
i01 i19 i26 i36
'4o3l}
-c se3 sa@ so0
i15 i34
uY2
i6c i7a
iB-stTi8z
sITD8+5
s2+ p4
-c so0 se3 sa@
'6i48
D4 $5
sr0
c ^f ^1
*96[
i1c i1d
-c sa@ so0 se3
i7a i3c
'6 Z1
c $0 $8 $7
.6 Z4
+6 s57
+3 +5 +8
-c so0 sa@ se3
Y3roA*
*30 i4l
Y5*94'A
*23x82
T8L5
i76 i57
+i6 c
D5 L6
f -6 x7B
+6 i4.
sa4 se3 si1 so0
i31 i59
o50 l
-sa4 si1 so0 se3
$Fo9]
i37 i47
sWB
o4e ^j
^b -3
s5@*26
-sa4 so0 se3 si1
+3 $4
'8 o5p K
i6k i7a i81
$5 *56
i3a i76
i51 $2
-oA[s2dsdX
+oA[s2X
o7ns8w$x
si1 T5
x53 i4l
-sa4 si1 se3 so0
$@$&
^D i86
*41 [ x21
o94 l
k o4k
.9sC_
-sa4 so0 si1 se3
l *20
x74 $y
$h xB3
i31 i52
^-xA6
^h ^s ^a
-sa4 se3 so0 si1
l $l $o $o $p
^o ,2
*31,7^D
i6h i7a
so%
i6d i7o
-so0 se3 sa@
E sLs
u d 'B $2
f x53
x42 i2h
i82 i90 iA2 iB3
k i3f
-c ss5 se3
$1 $9 $5 $3
x82 D5
Y3o7uo8R
+4 +4 -2
xA2 +A
i6Ks(HiBo
-so0 sa@ se3
+7 -5 -3
i4G
+6 $!
sfC
i4s i52
i0s i1o i2n
+i4 E
y1 f
s Zs@_
D6Y1
i17 i31
Z5sb9
^j+3
+i7 l
i8o i91
$s *21
^P $b
+3 }
u ^X ^x $X $x
$&$=
-s9Cs9W
.4xB8R4
x63 i4s
o4)
i31i49i57i67
^hx45
Y2 x32
-ss5 sa4 so0
k *31
c $0 $4 $3
*64 K
k i6c
]sKG$a
d o6y
-sDJ*7AsD@
+sDJ*7A
$A o00
x32 *46
'9 i2l
+0 -2 +6
^M z1
$ $l
+$ t
Z1 c +3
s4u
D8 *56
*04 +6
$x {
x04 $v
+i1 l
o7l -7 -3
-5 T2
R5 s1z
$e ^H
c o9!
+B x02
+o6 c
Z1 i2a
*51 ^o
*68 x52 {
*10 l
i6e i7r i8s
$A z1
-c sa@ sl1
i3a i4t i5a
Y1sFV
i73 i3f
i00 i18 i20
^2 ^G
^ ^3
+i8 {
*A1{sWY
^@ x15
[ c *31
*43 o13
xB2 $h
i1p i2o
+i5 t
xB9+A
$5 $5 $3 $1
*23 p1
+0 +6 +3
C*42
s9jY2K
+o4 t
$9 o84
c ^A ^A
+9$>,8
$9 $6 $! $!
x82 +7
y4 x46
-s7ys7)
+s7y
ssZy2
i67 i72
o9e K
t -7 *65
i9e iAn
R1x64
-so0 si! se3
l s1a
o86*53o9F
l i0M
-5 o0h
i0m i1u i2s
Y1 ^1
-c sa@ si!
c x53 Z3
x32 $8
c ^I
*A8 [
x16 i2n
i37 i73
-sDk*20sDF
+sDk*20
c ^f ^a
x56*34
o1G
x24 $8
^8iB`x46
i6a i7t i8e
-si1 sa@ se3
D5*5Ay1
c ^e ^3
x52 i01
Z2 x06
*B6 x62
*46 $v
-si1 se3 sa@
y3 p5 s9*
-1[o3Z
Z3x6B
o2r x43
$y o4f
D4i5DsP0
-c so0 si!
L7 l
c T3 i3_
*57K
o2m .6
sgCR6*61
^M *68
-so0 se3 si1 sa@
sbjt
sye
Z1 +7 -1
D6 i66
o7o o0d D8
i0l i1u i2i
-so0 si1 sa@ se3
+1 $0
-so0 sa@ se3 si1
D7 T6 *54
$s o1b
syc
D1 r {
x51 ^'
i6n i7g i81
-so0 se3 sa@ si1
i5i i6k
*17 *06
Z3 $8
l +1 o3w
i1V
x31 $M
-so0 sa@ si1 se3
*90p4T0
i5e i6m
i8u i9s
-4 ^F
T6 .1
o48 k
-so0 si1 se3 sa@
L6 $V
$. $p $k
Z1 i7n
i3S l r
i5f i70
^fry2
+o3 c
z5c
*54 k
i74 i1d
i1o i2m i3p
^R z1
^r stm
-c si! se3
x01Z5,9
y1 o6K
$. *43
D5 spz
$a *38
k ^8
-se3 ss$
x13 $g z1
i11 i73
x27 ^M
^W -6
^H *63
slx
-c si1 sa@ se3
s3q
T2 [ d
*51 D5
x54 $a
o2d x32
i0r ^z
-c se3 sa@ si1
x63 $a
q x48
i8v {
p3u
i6a z1
*78 i6k
-c se3 si1 sa@
$0 x26 Z2
L4o2Ik
o3}L5
o35 Z2
z3 o3j [
r +5
-o3U sUe ssa
.6sXv
c$1$2$3$<
c$1$2$3$>
i70 i5b
[ $z
*A7q
-: x03 l
+x03 l
x28*BAZ4
i5e i1d
i3- r }
] ] *13
Y2 y2
*05 *32
-s^6s6I
+s^I
D2 smP
T4 'A
scE
^_sV=
*65]
+3*1A
-c sa@ se3 si1
D1 $L
x78*67
xB5s|i^H
sse
kxB5sSO
iA9i6jiAf
-c ss$ si1
+1x7B
x03,B
$ s)3
$s $o $m $e
i5r i6d
i20 i30 i40
-c sa@ si1 se3
i5n i6y
i4t i5h i6a
i1o i2r i3n
xA1*B9
*08o5r*79
c ^9 $7
+$ $t
i30 i18
D0i9Vy1
[o7iy5
^X x13 +4
sfh
x74 i7|
-c ss$ so0
x19 Z4
sfz
y5 x42
*21y1x24
i5e R0 L0
c i1-
-si1 ss5
[ i0g
$,{x01
s7Vo4F
Z1 seu
r D4
ss9
+i1 C
ssk
ssT
Z2 x32
$9 $1 $@ $@
$8 $9 $* $*
$2 $9 $$ $$
-se3 si!
o66 $/
$9 $7 $$ $$
c $5 $3 $!
p2 'A
K s4G
oA2 xB2
-sa@ so0 si!
x47 Z4
*74 $k
$U .2
-0 -4 -8
D4 *87 l
+C +E +E
-c se3 ss$
Z2 $*
+6 +C
i2h }
+D -B -C
c T4 sa@
-siN sNk
+4 +8 +5
+E -C -C
i4k D5
{ Z2 +1
*69 o95
*54 o5l
-c ss$ sa@
*10 t x81
Z1 o39
Z2 x94
dY2
i3. .4
i50 T6
-so0 si1 sa@
x71 o4s
.6 $S
^m *42
x46 Z2 c
K *30
i31 ,5
-so0 sa@ si1
i03 i12
i03 i10
i00 z1
x42x57s*Q
i72 i80 oA7
Z5 l x95
-c se3 si1 so0
$L*47
R4 *46 Y1
'8Z5o81
T7o2J
i81 [
o92$l
-c si1 se3 so0
*74 k x24
o9jsOei1a
$Ni2R]
TB x61
p2 i4d
s$?
+i2 E
T3 d
i18 i1e
$gs^;
t x45 Z1
*10 $0
x72 [ l
-c se3 so0 si1
o2f s24
*40 ^G
z1 Z5
z5,4
suL
x72$z$k
-c so0 se3 si1
l '6 $7
i34 i5d
T4 TB
x93 *78
i0z u
K *B6
-so0 ss$ se3
x53x53*89
Cx56x87
{ x21 o5n
o4d *76
l $. $m $a
s9r
-c si1 so0 se3
o6k*8B
x57 i0a
s!(
^l{
p4 xA3 s0T
Z1 ^p
-se3 ss$ si1
i3e i35
y3 x54 *20
i04 ^!
sAjx0A
*BA c
$f $a $i $r $y
-se3 so0 sg9
sCO
i01sn`
u i26
x02 Z2 sBD
i54 x62
o60 *98
-se3 sg9 si1
o71 +8
K,7s?Y
i9hsiqo6U
*31 x29
*57R6o1k
i0b Z1
-sOusOG
+sOu
lsz<sLP
$0 $1 $# $#
^bx39
x57oBS
*4Bx73
x13 $h ^9
-D5 :
*91 *37
z2i9E
,6i5d
i16*85
i7v i8e i91
i79 i35
-D2 D4 :
R4 $9
c $1 $4 $@ $@
i64 *65
szj o54 s4y
o3i +2
T3 T8 T5 T7
-c sl1 si!
x05 o3t
$e *34 -2
Y1 x02 {
$@x5B*AB
i6m i8n i7a
i36 L2 L2
+i2 u
D0 +1
o5ro3Y
i3i i4l i5l
oAa x73 z1
sd}-5
T0 o4O
-se3 so0 sg9 si1
i6. o6o x72
o5o -6
^4 *70
s+Ii37L3
u ^2 ^b
-0 o2p
-se3 si1 so0 sg9
o62 .5
l $8 $9 $4
x52 o1P
D3s>h*87
i47 o1g
$e R3
-c si1 ss$ se3
i6bx4B^G
i0c i1o i2c
d r x0B
c $7 $8 $6 $!
*26 x15
*56 s06
+o2 k
*91*20x49
x02 o4#
s4y D5
l $. $m $r
R5 +2
k *54 R3
-c so0 si! se3
xA6 y2
s3`*53
c*B9x71
x71x48y2
sP1x42T9
x65 x03
-se3 so0 si1 sg9
T0*A1
$E x04 t
srS Y2
sJIi3C
^E o6b
.2 *41
-c se3 ss$ si1
c o9$
*26 x65
T0 T7 T6 T4 T5
l *50
i7v x68 *67
[ [ $2
-sfu sue
+sfe
o6L *75 t
i5r i60 i70
^R spm
s!E
x32 D1 $1
+8 R7
-c ss$ si1 se3
z4z3
x59 T1
L5 ^d
*43x29^X
^s ^a ^m
o7i *25 +3
-c si! se3 so0
i2u x72
sB8
^a ^t ^n ^i ^c
^+ ^7
c ^* $2
*40*97
-c ss$ se3 si1
c ^) $0
R5 st3
*30 x32 $t
R9C
y3z2l
'7 Y1
-se3 sg9 so0 si1
i70 i87
c ^4 $*
i0y $c
$3 u x42
*30 x35 sl$
x05'7
-sZEsEd
c ^1 ^-
c ^1 ^$
*30 $t x32
R6 x41
d 'E $9
[ o0m $i $a i4o
-c so0 se3 si!
*A2^vsMK
i4S x53
c ^1 $/
z3 x27
c ^! $3
$* ^9
-so0 si! ss$
T3 TA T7 T6
+1o5b
+5s1S,4
T0 T8 T3 T7
sWF,1*7A
D4 i4C
-c si1 se3 ss$
T1 T9 T2 T6
*83 x73
L5 i5h
$" u
T2 T7 T3 T4
$. $a $n
-sa4 sg9 se3
$N*58
i76 i3b
sCM
i44 $h
T0 T3 T9 T2
x71x53
-sc< si1 so0
spM o7t
i6n t
sw,i8$
*A5 *A6
^m ^i ^k
x29'B*3A
-c si! so0 se3
c $F $C
x46 -1
T2 T6 T3 T9
] R3 spk
sbF*A0'A
K i4+
-se3 si1 sg9 so0
s_I*25*93
*83R4
Es#U
.8x02
K i00
o6f z1
-c se3 si! so0
^x z3
T1 T3 T9
T0 d [
c ^n ^u
^k,7su|
o0s *24
+^ ^u
oBa'9x3B
L7s0q
x8B*45D1
s+qT9x81
o5w ^3
c ^y ^k ^s
-[ s8M soA
sO+o2YT1
x56*14'B
^" D4
x53 D1 *15
s'Y
oA)
-sXcs^Us^`
+sXcs^U
-C -C -E
D1 L3
i7f L6
$k $a $t $z
cux57
$b $a $r $u
+^ r
sT(sEz
x28i1J
$m $o $f $o
+B *78
z4 $?
'A o3D
-y2s|asaR
+y2s|R
.Ai4Z
x86x0A
x8A*1A
t ^I
slI
ss$ T8
-sgv sbg
L0 Z2 Z2
C i4i
} k u
$r $o $l $l
o7n t
$p $o $t $t $y
-sa@ si! se3
$s $n $o $w
$b $a $g $g $y
$l $o $r $d
^a ^t ^s ^i ^s
^d ^a ^a
^n ^i ^h ^s
-sa4 ss5
^h ^o ^o ^p
^i ^m ^i ^m
^a ^g ^o ^y
^a ^k ^e
^r ^o ^r ^r ^e
syV ^O
+i3 {
sc8
$0 d D4
sC* $T
o59 o31
p3$d
i7s y3
-sVusV*o1{
+sVuo1{
L6]sH&
o2k -5
Y5 Y2 x39
*56 i7,
x26 '6
o4m t
+oB c
x85 *53
{ { r
x31 ^x
] k ^o
i9c ] o7K
[ scR
+i4 r
^n +7
i19 R0 R2
saq y3
*47 $j
iB&x27*35
*47 $e
+i0 l
[ +2 o0t
i0c i1e
x34 som
Y3s=R
o5l '9
^k $z
-sg9 si1 so0
$4p5
o2@TA*B8
i8# *84 x42
*21kY5
xB4 oAX
s.}$,s+'
+o0 t
sr@
srE
srV
i8e iA0 i91
y3x5B
p5x3Ao5e
-se1 i32 sl7 st3 o75
[ x37 x41
*16x83*8A
r *21 *05
z5 x31 ^K
o1c ^p
.5 r i5c
-syMsyP
$^ Z2
d Y1 Z1
x3A f
u s6&
+4 *16
T1 T3 T8 TB
-sl1 sa@
o0[*60
l $. $s $m
Z4 x05
T0 T3 TB T2
snt r
-7 *64
-'5 p2
-l o81 RA
*53 ,2 *83
s41 *89
R2 *76
T1 so0
*52 *74
Y4 x29
+*73 scu
+*37 Y1
+s*A
o3n i2k
$p x76
-8 *07 K
T2 TD T7
*78 x92
$1 +4
-^9 o6f $d su3 se6
x26 f u
^| ^8
$1 +2
oAl x72
*53 o5z
c $5 $1 $7
-^4 i2b so2 sa3 o56
T1 TA T8 T7
sh. x93 r
*86 Y1
i5e L2 L7
i38 R7 L3
i44 i43 i42 i41
-+1 i67 sc5 ss0 o7a
D9 $1
i3a L5 L0
i71 L5 R1
^a ^m ^a ^k
*48 T7
x1A x5B
-x34 ssm som
i5a i6n i72
T1 T9 TA T2
} ^F c
u$6$f
C D5
sx|
-^0 $6 so0 sm0 ,6
C D4
u^b^c
u^f^A
l $8 $5 $3
l $8 $5 $2
x4A }
-su3 o79 sa0 sk8 o5k
T3 T7 TB T6
si$
u ^D ^7
T0 T5 T3 T9 T6
o6? *45
$d o4R
-^0 ,6 $6 so0 sm0
T0 T5 T2 T4 T8
slW
x63 Z3
s8$
saDxB7i5A
Z2 -4 i38
-^0 ,6 sm0 so0 $6
l $7 $2 $4
T0 T2 TA T8 T6
u $9 $0 $5
l $5 $7 $6
$^E
T1 T3 T5 T6 T2
+i4[
^F ^7
o9ss4%*19
T3 TB T6
l $9 $8 $6
{ x54
^1 *01 $0
-su3 sk8 sa0 o79 o5k
o0L p1
T0 T8 T5 T2 T4
T0 T8 T5 T2 T1
i96 i7s x37
l $4 $3 $7
T1 T2 TA T5
-^0 so0 sm0 ,6 $6
l $4 $3 $2
l $4 $3 $0
L2 i4e $-
K ^0 +0
$! ^/
k -1 o0k
-su3 sa0 sk8 o5k o79
$- $8
$+ ^#
x95s?[
$b $l $i $n $k
$b $l $i $n $d
o8$ E
-+1 ss0 sc5 i67 o7a
i9g iAo iBn
i5i i6e i71
^m ^a ^I
$f $i $l $e $s
i62 i80 i70 i96
o2#]
-c so0 sa@ si!
$h $a $l $e
^9 su3 o6f $d se6
p4 ,3
iBPsn=L9
i0? $!
d *64
-o1y $f sl1 se2 $l
-c si! so0 sa@
$c $a $n $s
$t $i $m $e $s
T1 T7 T8 T6
^Y z2
i89q
o1c o75 o3j o69 o5d
-c sa@ ss$ se3
$m $a $n $e
$g $r $o $w $l
^V +3
i44 i53 i62
$p $a $l $s
i2s i4r i3e
-c sa@ si! so0
$l $i $o $n $s
R3 E
Z3 sy7
T0 TB T3 T2
-3 ,2
-1 o26
-c ss$ se3 sa@
$b $i $r $d
i5u i6s i71
sy5
$s $t $e $w
$i $7
i6d x13
-c so0 si! sa@
c i11 i31 i51 i71
$i $r $o $n $s
c $. $m $t
$r $a $d $i $o
o87 xA2
T4 TA T7 T8
-c sa@ se3 ss$
o2>
u i11 i31 i51
Z1o1Lo1c
$Q Z1
x78 $j
$Y $2
-c ss$ si!
$d $e $l $t $a
i82 iA1 i92
i10 i20
$! i4- l
^0 ^k
T6 T8 T9 T7
-c si! sa@ so0
$d $o $d $g $e
x25 Z3
$c $u $b $e
-c se3 sa@ ss$
^o ^l ^b ^a ^p
i3t i5e i4l
i36 i45
$e $a $g $l $e
x29 K y4
$0 $_
-c se3 ss$ sa@
$i $t $e $m
$g $r $a $m
,6 i75 x41
^t ^e ^k
^o ^i ^l
i1h D4
-se1 i32 o75 sl7 st3
^m ^r ^e
i14 ss6 i35 $9 $3
i1w i2e i3r
^i ^d ^n ^a ^m
D5 i89 i97 o6g RA
^a ^d ^n ^a ^j
-^9 $d o6f su3 se6
^a ^d ^n ^a ^p
i4c i6t i5a
^Q u
^c T1
^i ^r ^o ^n
^t ^s ^u ^l
-^0 $6 ,6 so0 sm0
Z1 sjv
[ o0/
^y ^k ^c ^i ^r
T4 TA T8 TB
s5m [ *98
x12 ^z
-so0 ss5
T1 T2 T6 TA
^h ^t ^o ^g
^o ^b ^m ^a ^r
sY5 *52
i8# u
sy! [
-saG sGE
+saE
x13 x47
$I-9C
^d $4 ^g $u
x81 {
T0 TA T6 TE
y2*70$4
-si! so0
D3 sKs
*58 x05
d sC. x64
z3*A0x51
^n ^u ^a ^h ^s
sNj
-^9 se6 su3 $d o6f
} i31 +0
oA< x14
^u ^d ^a ^r
^a ^n ^i ^h
T6q
i8H x72 Y5
-se3 ss5
l q
^e ^d ^a ^h ^j
^v ^d ^a
^F*B0
$5 sLN
swe o3a
-+1 i67 o7a ss0 sc5
c $0 $4 $2 $!
^r ^e ^l ^y ^t
^u ^o ^n ^i ^m
$j $o $r $g $e
^g ^l ^t
c $0 $8 $6 $!
-sa@ ss5
o0e o4_ i8a ,9 o6a
c $0 $1 $9 $!
*1Bx85z3
^k ^s ^i ^d
[LB
o2j D6 o7d o34
-o1y sl1 se2 $f $l
T0 T8 TB T3
d i4b
x3B R2
*BAx01
$n $o $v $2 $1
TA^$,2
-^9 su3 se6 o6f $d
^n ^z ^a
l $. $w $s
+6 p4 $6
^8 x36
y2*36
i3y x51 Z2
-c sa4 si!
o3W i4A
T0 T5 TA T7
^c ^e ^l ^a
i6U x16
T2x96
o1co7$x49
+i3 c
^n ^a ^m ^b
.2 x53
^b ^m ^k
c $1 $4 $7 $!
$j $u $l $8
$r $u $n $e $s
+o3 }
-5 -7 -D
x63s5^
$t $i $t $o
svp D7 K
+D -E +F
R7*14
+s G u
+4 +7 +5
+C -6 +C
^o ^g ^a ^s
+0 -7 -2
x45 Y4
D3 x64
+i4{
+4 +8 -7
$!$`
^i ^a ^d
R2 x62
y1 *A9 +6
+0sz@
+s0K
x1A,9{
+9 $\
c $5 $6 $7 $!
*80 T5
T0 T7 TB TA
y2 x61 *05
-^9 $d su3 se6 o6f
D0 ^E
z3*02u
c $3 $4 $8 $!
x1B u i0A
sl x4B
c $9 $5 $7 $!
-se3 ss$ si!
si! ss$ u
i5t i7e i6t
-3 s0b
*19 x24
^\ z3
E^f
-^9 o6f su3 se6 $d
^so3Rx9A
c $6 $4 $3 $!
^p Z2
i4a k x71
i5o o2n
$f $l $y $s
-^4 so2 sa3 i2b o56
Z2 *90 $e
o1y sl1 $f se2 $l
$s $d $o $l
x94 x35
x43 i3t
^l ^u ^j c
-^0 so0 sm0 $6 ,6
$6 d D4
$6 d D7
$6 x43
i8@ x16
x53 $h
L3 z1
-se1 sl7 st3 i32 o75
[ r x51
+7 -A -6
s#p
Y1x38z2
o7k l D3
r s6k
-^4 i2b o56 sa3 so2
sI,
p2s u
x37 ^p
u ^7 ^6
D7i1Bs.|
i31 L3 L1
+i3 ]
Z1 i8y
Y2 *51 D3
i3Ri7'
*B3*4A
^h ^s ^a ^b
o2g k
+sg t
i44 Z1
i14 R7 L1
i34 R4 R1
sm> sa>
sms $
smr x32
-smR sRx $c
-smR $c sRx
+smx $c
smr
smq u
smq +2
$. sei
^s E D7
se_ d [
-sec y2 scp
+sep y2
sec Y1
-sec scp y2
sec ,3
se@ c
$S E $B
] sbt $7
sbt ,3
sbS Y2
-sbs ssK o5*
-sbs o5* ssK
+sbK o5*
sb sBO
sbs *25
sbs
^S *21 R2
s$2 ^1
} s21
-s20 s0@
^s *20 D4
s20 +5
$S *20 *35
o0a i3e
i47 i58 o66
+3 $o
-se3 so0
o4a o5z r
$h $1 $2 $3
o2k ,3
#define MARKOV_DISABLE 0
#define MARKOV_CLASSIC 0
#define BENCHMARK 0
-#define BENCHMARK_REPEATS 2
+#define BENCHMARK_REPEATS 100
#define RESTORE 0
#define RESTORE_TIMER 60
#define RESTORE_DISABLE 0
if (device_param->skipped) continue;
- u64 speed_cnt = 0;
- float speed_ms = 0;
+ u64 speed_cnt = 0;
+ double speed_ms = 0;
for (int i = 0; i < SPEED_CACHE; i++)
{
- float rec_ms;
+ double rec_ms;
hc_timer_get (device_param->speed_rec[i], rec_ms);
* speed new
*/
- u64 speed_cnt[DEVICES_MAX] = { 0 };
- float speed_ms[DEVICES_MAX] = { 0 };
+ u64 speed_cnt[DEVICES_MAX] = { 0 };
+ double speed_ms[DEVICES_MAX] = { 0 };
for (uint device_id = 0; device_id < data.devices_cnt; device_id++)
{
for (int i = 0; i < SPEED_CACHE; i++)
{
- float rec_ms;
+ double rec_ms;
hc_timer_get (device_param->speed_rec[i], rec_ms);
* timers
*/
- float ms_running = 0;
+ double ms_running = 0;
hc_timer_get (data.timer_running, ms_running);
- float ms_paused = data.ms_paused;
+ double ms_paused = data.ms_paused;
if (data.devices_status == STATUS_PAUSED)
{
- float ms_paused_tmp = 0;
+ double ms_paused_tmp = 0;
hc_timer_get (data.timer_paused, ms_paused_tmp);
}
}
- float ms_real = ms_running - ms_paused;
+ double ms_real = ms_running - ms_paused;
float cpt_avg_min = (float) data.cpt_total / ((ms_real / 1000) / 60);
float cpt_avg_hour = (float) data.cpt_total / ((ms_real / 1000) / 3600);
static void status_benchmark ()
{
- if (data.devices_status == STATUS_INIT) return;
+ if (data.devices_status == STATUS_INIT) return;
if (data.devices_status == STATUS_STARTING) return;
if (data.words_cnt == 0) return;
- u64 speed_cnt[DEVICES_MAX] = { 0 };
- float speed_ms[DEVICES_MAX] = { 0 };
+ u64 speed_cnt[DEVICES_MAX] = { 0 };
+ double speed_ms[DEVICES_MAX] = { 0 };
for (uint device_id = 0; device_id < data.devices_cnt; device_id++)
{
if (device_param->skipped) continue;
- speed_cnt[device_id] = 0;
- speed_ms[device_id] = 0;
-
- for (int i = 0; i < SPEED_CACHE; i++)
- {
- speed_cnt[device_id] += device_param->speed_cnt[i];
- speed_ms[device_id] += device_param->speed_ms[i];
- }
-
- speed_cnt[device_id] /= SPEED_CACHE;
- speed_ms[device_id] /= SPEED_CACHE;
+ speed_cnt[device_id] = device_param->speed_cnt[0];
+ speed_ms[device_id] = device_param->speed_ms[0];
}
float hashes_all_ms = 0;
if (event_update)
{
- float exec_time;
+ double exec_time;
hc_timer_get (timer, exec_time);
if (data.benchmark == 1)
{
- for (u32 i = 0; i < data.benchmark_repeats; i++)
+ double exec_ms_avg_prev = get_avg_exec_time (device_param, EXEC_CACHE);
+
+ // a few caching rounds
+
+ for (u32 i = 0; i < 2; i++)
{
+ hc_timer_set (&device_param->timer_speed);
+
choose_kernel (device_param, data.attack_exec, data.attack_mode, data.opts_type, salt_buf, highest_pw_len, pws_cnt);
+
+ double exec_ms_avg = get_avg_exec_time (device_param, EXEC_CACHE);
+
+ exec_ms_avg_prev = exec_ms_avg;
+ }
+
+ // benchmark_repeats became a maximum possible repeats
+
+ for (u32 i = 2; i < data.benchmark_repeats; i++)
+ {
+ hc_timer_set (&device_param->timer_speed);
+
+ choose_kernel (device_param, data.attack_exec, data.attack_mode, data.opts_type, salt_buf, highest_pw_len, pws_cnt);
+
+ double exec_ms_avg = get_avg_exec_time (device_param, EXEC_CACHE);
+
+ if ((exec_ms_avg_prev / exec_ms_avg) < 1.001) break;
+
+ exec_ms_avg_prev = exec_ms_avg;
}
}
u64 perf_sum_all = (u64) pws_cnt * (u64) innerloop_left;
- if (data.benchmark == 1)
- {
- perf_sum_all = (perf_sum_all * data.benchmark_repeats) + perf_sum_all;
- }
-
hc_thread_mutex_lock (mux_counter);
data.words_progress_done[salt_pos] += perf_sum_all;
* speed
*/
- float speed_ms;
+ double speed_ms;
hc_timer_get (device_param->timer_speed, speed_ms);
hc_thread_mutex_lock (mux_display);
+ // current speed
+
device_param->speed_cnt[speed_pos] = perf_sum_all;
device_param->speed_ms[speed_pos] = speed_ms;
speed_pos = 0;
}
+ // average speed
+
+ device_param->speed_cnt_total += perf_sum_all;
+
+ device_param->speed_ms_total += speed_ms;
+
/**
* benchmark
*/
if (data.quiet == 0) log_info ("");
/**
- * Inform user which algorithm is checked and at which workload setting
+ * In benchmark-mode, inform user which algorithm is checked
*/
if (benchmark == 1)
device_param->speed_pos = 0;
memset (device_param->speed_cnt, 0, SPEED_CACHE * sizeof (u64));
- memset (device_param->speed_ms, 0, SPEED_CACHE * sizeof (float));
+ memset (device_param->speed_ms, 0, SPEED_CACHE * sizeof (double));
memset (device_param->speed_rec, 0, SPEED_CACHE * sizeof (hc_timer_t));
+ device_param->speed_cnt_total = 0;
+ device_param->speed_ms_total = 0;
+
device_param->exec_pos = 0;
memset (device_param->exec_ms, 0, EXEC_CACHE * sizeof (double));
{
if (data.devices_status == STATUS_PAUSED)
{
- float ms_paused;
+ double ms_paused;
hc_timer_get (data.timer_paused, ms_paused);
projects / hashcat.git / commitdiff