#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) ? -1 : 0)
+ #define K01 (((w0s >> ( 0 + 6)) & 1) ? -1 : 0)
+ #define K02 (((w0s >> ( 0 + 5)) & 1) ? -1 : 0)
+ #define K03 (((w0s >> ( 0 + 4)) & 1) ? -1 : 0)
+ #define K04 (((w0s >> ( 0 + 3)) & 1) ? -1 : 0)
+ #define K05 (((w0s >> ( 0 + 2)) & 1) ? -1 : 0)
+ #define K06 (((w0s >> ( 0 + 1)) & 1) ? -1 : 0)
+ #define K07 (((w0s >> ( 8 + 7)) & 1) ? -1 : 0)
+ #define K08 (((w0s >> ( 8 + 6)) & 1) ? -1 : 0)
+ #define K09 (((w0s >> ( 8 + 5)) & 1) ? -1 : 0)
+ #define K10 (((w0s >> ( 8 + 4)) & 1) ? -1 : 0)
+ #define K11 (((w0s >> ( 8 + 3)) & 1) ? -1 : 0)
+ #define K12 (((w0s >> ( 8 + 2)) & 1) ? -1 : 0)
+ #define K13 (((w0s >> ( 8 + 1)) & 1) ? -1 : 0)
+ #define K14 (((w0s >> (16 + 7)) & 1) ? -1 : 0)
+ #define K15 (((w0s >> (16 + 6)) & 1) ? -1 : 0)
+ #define K16 (((w0s >> (16 + 5)) & 1) ? -1 : 0)
+ #define K17 (((w0s >> (16 + 4)) & 1) ? -1 : 0)
+ #define K18 (((w0s >> (16 + 3)) & 1) ? -1 : 0)
+ #define K19 (((w0s >> (16 + 2)) & 1) ? -1 : 0)
+ #define K20 (((w0s >> (16 + 1)) & 1) ? -1 : 0)
+ #define K21 (((w0s >> (24 + 7)) & 1) ? -1 : 0)
+ #define K22 (((w0s >> (24 + 6)) & 1) ? -1 : 0)
+ #define K23 (((w0s >> (24 + 5)) & 1) ? -1 : 0)
+ #define K24 (((w0s >> (24 + 4)) & 1) ? -1 : 0)
+ #define K25 (((w0s >> (24 + 3)) & 1) ? -1 : 0)
+ #define K26 (((w0s >> (24 + 2)) & 1) ? -1 : 0)
+ #define K27 (((w0s >> (24 + 1)) & 1) ? -1 : 0)
+ #define K28 (((w1s >> ( 0 + 7)) & 1) ? -1 : 0)
+ #define K29 (((w1s >> ( 0 + 6)) & 1) ? -1 : 0)
+ #define K30 (((w1s >> ( 0 + 5)) & 1) ? -1 : 0)
+ #define K31 (((w1s >> ( 0 + 4)) & 1) ? -1 : 0)
+ #define K32 (((w1s >> ( 0 + 3)) & 1) ? -1 : 0)
+ #define K33 (((w1s >> ( 0 + 2)) & 1) ? -1 : 0)
+ #define K34 (((w1s >> ( 0 + 1)) & 1) ? -1 : 0)
+ #define K35 (((w1s >> ( 8 + 7)) & 1) ? -1 : 0)
+ #define K36 (((w1s >> ( 8 + 6)) & 1) ? -1 : 0)
+ #define K37 (((w1s >> ( 8 + 5)) & 1) ? -1 : 0)
+ #define K38 (((w1s >> ( 8 + 4)) & 1) ? -1 : 0)
+ #define K39 (((w1s >> ( 8 + 3)) & 1) ? -1 : 0)
+ #define K40 (((w1s >> ( 8 + 2)) & 1) ? -1 : 0)
+ #define K41 (((w1s >> ( 8 + 1)) & 1) ? -1 : 0)
+ #define K42 (((w1s >> (16 + 7)) & 1) ? -1 : 0)
+ #define K43 (((w1s >> (16 + 6)) & 1) ? -1 : 0)
+ #define K44 (((w1s >> (16 + 5)) & 1) ? -1 : 0)
+ #define K45 (((w1s >> (16 + 4)) & 1) ? -1 : 0)
+ #define K46 (((w1s >> (16 + 3)) & 1) ? -1 : 0)
+ #define K47 (((w1s >> (16 + 2)) & 1) ? -1 : 0)
+ #define K48 (((w1s >> (16 + 1)) & 1) ? -1 : 0)
+ #define K49 (((w1s >> (24 + 7)) & 1) ? -1 : 0)
+ #define K50 (((w1s >> (24 + 6)) & 1) ? -1 : 0)
+ #define K51 (((w1s >> (24 + 5)) & 1) ? -1 : 0)
+ #define K52 (((w1s >> (24 + 4)) & 1) ? -1 : 0)
+ #define K53 (((w1s >> (24 + 3)) & 1) ? -1 : 0)
+ #define K54 (((w1s >> (24 + 2)) & 1) ? -1 : 0)
+ #define K55 (((w1s >> (24 + 1)) & 1) ? -1 : 0)
/**
* 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) ? -1 : 0)
+ #define S01 (((s0 >> 1) & 1) ? -1 : 0)
+ #define S02 (((s0 >> 2) & 1) ? -1 : 0)
+ #define S03 (((s0 >> 3) & 1) ? -1 : 0)
+ #define S04 (((s0 >> 4) & 1) ? -1 : 0)
+ #define S05 (((s0 >> 5) & 1) ? -1 : 0)
+ #define S06 (((s0 >> 6) & 1) ? -1 : 0)
+ #define S07 (((s0 >> 7) & 1) ? -1 : 0)
+ #define S08 (((s0 >> 8) & 1) ? -1 : 0)
+ #define S09 (((s0 >> 9) & 1) ? -1 : 0)
+ #define S10 (((s0 >> 10) & 1) ? -1 : 0)
+ #define S11 (((s0 >> 11) & 1) ? -1 : 0)
+ #define S12 (((s0 >> 12) & 1) ? -1 : 0)
+ #define S13 (((s0 >> 13) & 1) ? -1 : 0)
+ #define S14 (((s0 >> 14) & 1) ? -1 : 0)
+ #define S15 (((s0 >> 15) & 1) ? -1 : 0)
+ #define S16 (((s0 >> 16) & 1) ? -1 : 0)
+ #define S17 (((s0 >> 17) & 1) ? -1 : 0)
+ #define S18 (((s0 >> 18) & 1) ? -1 : 0)
+ #define S19 (((s0 >> 19) & 1) ? -1 : 0)
+ #define S20 (((s0 >> 20) & 1) ? -1 : 0)
+ #define S21 (((s0 >> 21) & 1) ? -1 : 0)
+ #define S22 (((s0 >> 22) & 1) ? -1 : 0)
+ #define S23 (((s0 >> 23) & 1) ? -1 : 0)
+ #define S24 (((s0 >> 24) & 1) ? -1 : 0)
+ #define S25 (((s0 >> 25) & 1) ? -1 : 0)
+ #define S26 (((s0 >> 26) & 1) ? -1 : 0)
+ #define S27 (((s0 >> 27) & 1) ? -1 : 0)
+ #define S28 (((s0 >> 28) & 1) ? -1 : 0)
+ #define S29 (((s0 >> 29) & 1) ? -1 : 0)
+ #define S30 (((s0 >> 30) & 1) ? -1 : 0)
+ #define S31 (((s0 >> 31) & 1) ? -1 : 0)
+ #define S32 (((s1 >> 0) & 1) ? -1 : 0)
+ #define S33 (((s1 >> 1) & 1) ? -1 : 0)
+ #define S34 (((s1 >> 2) & 1) ? -1 : 0)
+ #define S35 (((s1 >> 3) & 1) ? -1 : 0)
+ #define S36 (((s1 >> 4) & 1) ? -1 : 0)
+ #define S37 (((s1 >> 5) & 1) ? -1 : 0)
+ #define S38 (((s1 >> 6) & 1) ? -1 : 0)
+ #define S39 (((s1 >> 7) & 1) ? -1 : 0)
+ #define S40 (((s1 >> 8) & 1) ? -1 : 0)
+ #define S41 (((s1 >> 9) & 1) ? -1 : 0)
+ #define S42 (((s1 >> 10) & 1) ? -1 : 0)
+ #define S43 (((s1 >> 11) & 1) ? -1 : 0)
+ #define S44 (((s1 >> 12) & 1) ? -1 : 0)
+ #define S45 (((s1 >> 13) & 1) ? -1 : 0)
+ #define S46 (((s1 >> 14) & 1) ? -1 : 0)
+ #define S47 (((s1 >> 15) & 1) ? -1 : 0)
+ #define S48 (((s1 >> 16) & 1) ? -1 : 0)
+ #define S49 (((s1 >> 17) & 1) ? -1 : 0)
+ #define S50 (((s1 >> 18) & 1) ? -1 : 0)
+ #define S51 (((s1 >> 19) & 1) ? -1 : 0)
+ #define S52 (((s1 >> 20) & 1) ? -1 : 0)
+ #define S53 (((s1 >> 21) & 1) ? -1 : 0)
+ #define S54 (((s1 >> 22) & 1) ? -1 : 0)
+ #define S55 (((s1 >> 23) & 1) ? -1 : 0)
+ #define S56 (((s1 >> 24) & 1) ? -1 : 0)
+ #define S57 (((s1 >> 25) & 1) ? -1 : 0)
+ #define S58 (((s1 >> 26) & 1) ? -1 : 0)
+ #define S59 (((s1 >> 27) & 1) ? -1 : 0)
+ #define S60 (((s1 >> 28) & 1) ? -1 : 0)
+ #define S61 (((s1 >> 29) & 1) ? -1 : 0)
+ #define S62 (((s1 >> 30) & 1) ? -1 : 0)
+ #define S63 (((s1 >> 31) & 1) ? -1 : 0)
/**
* 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) ? -1 : 0)
+ #define K01 (((w0s >> ( 0 + 6)) & 1) ? -1 : 0)
+ #define K02 (((w0s >> ( 0 + 5)) & 1) ? -1 : 0)
+ #define K03 (((w0s >> ( 0 + 4)) & 1) ? -1 : 0)
+ #define K04 (((w0s >> ( 0 + 3)) & 1) ? -1 : 0)
+ #define K05 (((w0s >> ( 0 + 2)) & 1) ? -1 : 0)
+ #define K06 (((w0s >> ( 0 + 1)) & 1) ? -1 : 0)
+ #define K07 (((w0s >> ( 8 + 7)) & 1) ? -1 : 0)
+ #define K08 (((w0s >> ( 8 + 6)) & 1) ? -1 : 0)
+ #define K09 (((w0s >> ( 8 + 5)) & 1) ? -1 : 0)
+ #define K10 (((w0s >> ( 8 + 4)) & 1) ? -1 : 0)
+ #define K11 (((w0s >> ( 8 + 3)) & 1) ? -1 : 0)
+ #define K12 (((w0s >> ( 8 + 2)) & 1) ? -1 : 0)
+ #define K13 (((w0s >> ( 8 + 1)) & 1) ? -1 : 0)
+ #define K14 (((w0s >> (16 + 7)) & 1) ? -1 : 0)
+ #define K15 (((w0s >> (16 + 6)) & 1) ? -1 : 0)
+ #define K16 (((w0s >> (16 + 5)) & 1) ? -1 : 0)
+ #define K17 (((w0s >> (16 + 4)) & 1) ? -1 : 0)
+ #define K18 (((w0s >> (16 + 3)) & 1) ? -1 : 0)
+ #define K19 (((w0s >> (16 + 2)) & 1) ? -1 : 0)
+ #define K20 (((w0s >> (16 + 1)) & 1) ? -1 : 0)
+ #define K21 (((w0s >> (24 + 7)) & 1) ? -1 : 0)
+ #define K22 (((w0s >> (24 + 6)) & 1) ? -1 : 0)
+ #define K23 (((w0s >> (24 + 5)) & 1) ? -1 : 0)
+ #define K24 (((w0s >> (24 + 4)) & 1) ? -1 : 0)
+ #define K25 (((w0s >> (24 + 3)) & 1) ? -1 : 0)
+ #define K26 (((w0s >> (24 + 2)) & 1) ? -1 : 0)
+ #define K27 (((w0s >> (24 + 1)) & 1) ? -1 : 0)
+ #define K28 (((w1s >> ( 0 + 7)) & 1) ? -1 : 0)
+ #define K29 (((w1s >> ( 0 + 6)) & 1) ? -1 : 0)
+ #define K30 (((w1s >> ( 0 + 5)) & 1) ? -1 : 0)
+ #define K31 (((w1s >> ( 0 + 4)) & 1) ? -1 : 0)
+ #define K32 (((w1s >> ( 0 + 3)) & 1) ? -1 : 0)
+ #define K33 (((w1s >> ( 0 + 2)) & 1) ? -1 : 0)
+ #define K34 (((w1s >> ( 0 + 1)) & 1) ? -1 : 0)
+ #define K35 (((w1s >> ( 8 + 7)) & 1) ? -1 : 0)
+ #define K36 (((w1s >> ( 8 + 6)) & 1) ? -1 : 0)
+ #define K37 (((w1s >> ( 8 + 5)) & 1) ? -1 : 0)
+ #define K38 (((w1s >> ( 8 + 4)) & 1) ? -1 : 0)
+ #define K39 (((w1s >> ( 8 + 3)) & 1) ? -1 : 0)
+ #define K40 (((w1s >> ( 8 + 2)) & 1) ? -1 : 0)
+ #define K41 (((w1s >> ( 8 + 1)) & 1) ? -1 : 0)
+ #define K42 (((w1s >> (16 + 7)) & 1) ? -1 : 0)
+ #define K43 (((w1s >> (16 + 6)) & 1) ? -1 : 0)
+ #define K44 (((w1s >> (16 + 5)) & 1) ? -1 : 0)
+ #define K45 (((w1s >> (16 + 4)) & 1) ? -1 : 0)
+ #define K46 (((w1s >> (16 + 3)) & 1) ? -1 : 0)
+ #define K47 (((w1s >> (16 + 2)) & 1) ? -1 : 0)
+ #define K48 (((w1s >> (16 + 1)) & 1) ? -1 : 0)
+ #define K49 (((w1s >> (24 + 7)) & 1) ? -1 : 0)
+ #define K50 (((w1s >> (24 + 6)) & 1) ? -1 : 0)
+ #define K51 (((w1s >> (24 + 5)) & 1) ? -1 : 0)
+ #define K52 (((w1s >> (24 + 4)) & 1) ? -1 : 0)
+ #define K53 (((w1s >> (24 + 3)) & 1) ? -1 : 0)
+ #define K54 (((w1s >> (24 + 2)) & 1) ? -1 : 0)
+ #define K55 (((w1s >> (24 + 1)) & 1) ? -1 : 0)
/**
* 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) ? -1 : 0)
+ #define K01 (((w0s >> ( 0 + 6)) & 1) ? -1 : 0)
+ #define K02 (((w0s >> ( 0 + 5)) & 1) ? -1 : 0)
+ #define K03 (((w0s >> ( 0 + 4)) & 1) ? -1 : 0)
+ #define K04 (((w0s >> ( 0 + 3)) & 1) ? -1 : 0)
+ #define K05 (((w0s >> ( 0 + 2)) & 1) ? -1 : 0)
+ #define K06 (((w0s >> ( 0 + 1)) & 1) ? -1 : 0)
+ #define K07 (((w0s >> ( 0 + 0)) & 1) ? -1 : 0)
+ #define K08 (((w0s >> ( 8 + 7)) & 1) ? -1 : 0)
+ #define K09 (((w0s >> ( 8 + 6)) & 1) ? -1 : 0)
+ #define K10 (((w0s >> ( 8 + 5)) & 1) ? -1 : 0)
+ #define K11 (((w0s >> ( 8 + 4)) & 1) ? -1 : 0)
+ #define K12 (((w0s >> ( 8 + 3)) & 1) ? -1 : 0)
+ #define K13 (((w0s >> ( 8 + 2)) & 1) ? -1 : 0)
+ #define K14 (((w0s >> ( 8 + 1)) & 1) ? -1 : 0)
+ #define K15 (((w0s >> ( 8 + 0)) & 1) ? -1 : 0)
+ #define K16 (((w0s >> (16 + 7)) & 1) ? -1 : 0)
+ #define K17 (((w0s >> (16 + 6)) & 1) ? -1 : 0)
+ #define K18 (((w0s >> (16 + 5)) & 1) ? -1 : 0)
+ #define K19 (((w0s >> (16 + 4)) & 1) ? -1 : 0)
+ #define K20 (((w0s >> (16 + 3)) & 1) ? -1 : 0)
+ #define K21 (((w0s >> (16 + 2)) & 1) ? -1 : 0)
+ #define K22 (((w0s >> (16 + 1)) & 1) ? -1 : 0)
+ #define K23 (((w0s >> (16 + 0)) & 1) ? -1 : 0)
+ #define K24 (((w0s >> (24 + 7)) & 1) ? -1 : 0)
+ #define K25 (((w0s >> (24 + 6)) & 1) ? -1 : 0)
+ #define K26 (((w0s >> (24 + 5)) & 1) ? -1 : 0)
+ #define K27 (((w0s >> (24 + 4)) & 1) ? -1 : 0)
+ #define K28 (((w0s >> (24 + 3)) & 1) ? -1 : 0)
+ #define K29 (((w0s >> (24 + 2)) & 1) ? -1 : 0)
+ #define K30 (((w0s >> (24 + 1)) & 1) ? -1 : 0)
+ #define K31 (((w0s >> (24 + 0)) & 1) ? -1 : 0)
+ #define K32 (((w1s >> ( 0 + 7)) & 1) ? -1 : 0)
+ #define K33 (((w1s >> ( 0 + 6)) & 1) ? -1 : 0)
+ #define K34 (((w1s >> ( 0 + 5)) & 1) ? -1 : 0)
+ #define K35 (((w1s >> ( 0 + 4)) & 1) ? -1 : 0)
+ #define K36 (((w1s >> ( 0 + 3)) & 1) ? -1 : 0)
+ #define K37 (((w1s >> ( 0 + 2)) & 1) ? -1 : 0)
+ #define K38 (((w1s >> ( 0 + 1)) & 1) ? -1 : 0)
+ #define K39 (((w1s >> ( 0 + 0)) & 1) ? -1 : 0)
+ #define K40 (((w1s >> ( 8 + 7)) & 1) ? -1 : 0)
+ #define K41 (((w1s >> ( 8 + 6)) & 1) ? -1 : 0)
+ #define K42 (((w1s >> ( 8 + 5)) & 1) ? -1 : 0)
+ #define K43 (((w1s >> ( 8 + 4)) & 1) ? -1 : 0)
+ #define K44 (((w1s >> ( 8 + 3)) & 1) ? -1 : 0)
+ #define K45 (((w1s >> ( 8 + 2)) & 1) ? -1 : 0)
+ #define K46 (((w1s >> ( 8 + 1)) & 1) ? -1 : 0)
+ #define K47 (((w1s >> ( 8 + 0)) & 1) ? -1 : 0)
+ #define K48 (((w1s >> (16 + 7)) & 1) ? -1 : 0)
+ #define K49 (((w1s >> (16 + 6)) & 1) ? -1 : 0)
+ #define K50 (((w1s >> (16 + 5)) & 1) ? -1 : 0)
+ #define K51 (((w1s >> (16 + 4)) & 1) ? -1 : 0)
+ #define K52 (((w1s >> (16 + 3)) & 1) ? -1 : 0)
+ #define K53 (((w1s >> (16 + 2)) & 1) ? -1 : 0)
+ #define K54 (((w1s >> (16 + 1)) & 1) ? -1 : 0)
+ #define K55 (((w1s >> (16 + 0)) & 1) ? -1 : 0)
/**
* 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) ? -1 : 0)
+ #define S01 (((s0 >> 1) & 1) ? -1 : 0)
+ #define S02 (((s0 >> 2) & 1) ? -1 : 0)
+ #define S03 (((s0 >> 3) & 1) ? -1 : 0)
+ #define S04 (((s0 >> 4) & 1) ? -1 : 0)
+ #define S05 (((s0 >> 5) & 1) ? -1 : 0)
+ #define S06 (((s0 >> 6) & 1) ? -1 : 0)
+ #define S07 (((s0 >> 7) & 1) ? -1 : 0)
+ #define S08 (((s0 >> 8) & 1) ? -1 : 0)
+ #define S09 (((s0 >> 9) & 1) ? -1 : 0)
+ #define S10 (((s0 >> 10) & 1) ? -1 : 0)
+ #define S11 (((s0 >> 11) & 1) ? -1 : 0)
+ #define S12 (((s0 >> 12) & 1) ? -1 : 0)
+ #define S13 (((s0 >> 13) & 1) ? -1 : 0)
+ #define S14 (((s0 >> 14) & 1) ? -1 : 0)
+ #define S15 (((s0 >> 15) & 1) ? -1 : 0)
+ #define S16 (((s0 >> 16) & 1) ? -1 : 0)
+ #define S17 (((s0 >> 17) & 1) ? -1 : 0)
+ #define S18 (((s0 >> 18) & 1) ? -1 : 0)
+ #define S19 (((s0 >> 19) & 1) ? -1 : 0)
+ #define S20 (((s0 >> 20) & 1) ? -1 : 0)
+ #define S21 (((s0 >> 21) & 1) ? -1 : 0)
+ #define S22 (((s0 >> 22) & 1) ? -1 : 0)
+ #define S23 (((s0 >> 23) & 1) ? -1 : 0)
+ #define S24 (((s0 >> 24) & 1) ? -1 : 0)
+ #define S25 (((s0 >> 25) & 1) ? -1 : 0)
+ #define S26 (((s0 >> 26) & 1) ? -1 : 0)
+ #define S27 (((s0 >> 27) & 1) ? -1 : 0)
+ #define S28 (((s0 >> 28) & 1) ? -1 : 0)
+ #define S29 (((s0 >> 29) & 1) ? -1 : 0)
+ #define S30 (((s0 >> 30) & 1) ? -1 : 0)
+ #define S31 (((s0 >> 31) & 1) ? -1 : 0)
+ #define S32 (((s1 >> 0) & 1) ? -1 : 0)
+ #define S33 (((s1 >> 1) & 1) ? -1 : 0)
+ #define S34 (((s1 >> 2) & 1) ? -1 : 0)
+ #define S35 (((s1 >> 3) & 1) ? -1 : 0)
+ #define S36 (((s1 >> 4) & 1) ? -1 : 0)
+ #define S37 (((s1 >> 5) & 1) ? -1 : 0)
+ #define S38 (((s1 >> 6) & 1) ? -1 : 0)
+ #define S39 (((s1 >> 7) & 1) ? -1 : 0)
+ #define S40 (((s1 >> 8) & 1) ? -1 : 0)
+ #define S41 (((s1 >> 9) & 1) ? -1 : 0)
+ #define S42 (((s1 >> 10) & 1) ? -1 : 0)
+ #define S43 (((s1 >> 11) & 1) ? -1 : 0)
+ #define S44 (((s1 >> 12) & 1) ? -1 : 0)
+ #define S45 (((s1 >> 13) & 1) ? -1 : 0)
+ #define S46 (((s1 >> 14) & 1) ? -1 : 0)
+ #define S47 (((s1 >> 15) & 1) ? -1 : 0)
+ #define S48 (((s1 >> 16) & 1) ? -1 : 0)
+ #define S49 (((s1 >> 17) & 1) ? -1 : 0)
+ #define S50 (((s1 >> 18) & 1) ? -1 : 0)
+ #define S51 (((s1 >> 19) & 1) ? -1 : 0)
+ #define S52 (((s1 >> 20) & 1) ? -1 : 0)
+ #define S53 (((s1 >> 21) & 1) ? -1 : 0)
+ #define S54 (((s1 >> 22) & 1) ? -1 : 0)
+ #define S55 (((s1 >> 23) & 1) ? -1 : 0)
+ #define S56 (((s1 >> 24) & 1) ? -1 : 0)
+ #define S57 (((s1 >> 25) & 1) ? -1 : 0)
+ #define S58 (((s1 >> 26) & 1) ? -1 : 0)
+ #define S59 (((s1 >> 27) & 1) ? -1 : 0)
+ #define S60 (((s1 >> 28) & 1) ? -1 : 0)
+ #define S61 (((s1 >> 29) & 1) ? -1 : 0)
+ #define S62 (((s1 >> 30) & 1) ? -1 : 0)
+ #define S63 (((s1 >> 31) & 1) ? -1 : 0)
/**
* 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) ? -1 : 0)
+ #define K01 (((w0s >> ( 0 + 6)) & 1) ? -1 : 0)
+ #define K02 (((w0s >> ( 0 + 5)) & 1) ? -1 : 0)
+ #define K03 (((w0s >> ( 0 + 4)) & 1) ? -1 : 0)
+ #define K04 (((w0s >> ( 0 + 3)) & 1) ? -1 : 0)
+ #define K05 (((w0s >> ( 0 + 2)) & 1) ? -1 : 0)
+ #define K06 (((w0s >> ( 0 + 1)) & 1) ? -1 : 0)
+ #define K07 (((w0s >> ( 0 + 0)) & 1) ? -1 : 0)
+ #define K08 (((w0s >> ( 8 + 7)) & 1) ? -1 : 0)
+ #define K09 (((w0s >> ( 8 + 6)) & 1) ? -1 : 0)
+ #define K10 (((w0s >> ( 8 + 5)) & 1) ? -1 : 0)
+ #define K11 (((w0s >> ( 8 + 4)) & 1) ? -1 : 0)
+ #define K12 (((w0s >> ( 8 + 3)) & 1) ? -1 : 0)
+ #define K13 (((w0s >> ( 8 + 2)) & 1) ? -1 : 0)
+ #define K14 (((w0s >> ( 8 + 1)) & 1) ? -1 : 0)
+ #define K15 (((w0s >> ( 8 + 0)) & 1) ? -1 : 0)
+ #define K16 (((w0s >> (16 + 7)) & 1) ? -1 : 0)
+ #define K17 (((w0s >> (16 + 6)) & 1) ? -1 : 0)
+ #define K18 (((w0s >> (16 + 5)) & 1) ? -1 : 0)
+ #define K19 (((w0s >> (16 + 4)) & 1) ? -1 : 0)
+ #define K20 (((w0s >> (16 + 3)) & 1) ? -1 : 0)
+ #define K21 (((w0s >> (16 + 2)) & 1) ? -1 : 0)
+ #define K22 (((w0s >> (16 + 1)) & 1) ? -1 : 0)
+ #define K23 (((w0s >> (16 + 0)) & 1) ? -1 : 0)
+ #define K24 (((w0s >> (24 + 7)) & 1) ? -1 : 0)
+ #define K25 (((w0s >> (24 + 6)) & 1) ? -1 : 0)
+ #define K26 (((w0s >> (24 + 5)) & 1) ? -1 : 0)
+ #define K27 (((w0s >> (24 + 4)) & 1) ? -1 : 0)
+ #define K28 (((w0s >> (24 + 3)) & 1) ? -1 : 0)
+ #define K29 (((w0s >> (24 + 2)) & 1) ? -1 : 0)
+ #define K30 (((w0s >> (24 + 1)) & 1) ? -1 : 0)
+ #define K31 (((w0s >> (24 + 0)) & 1) ? -1 : 0)
+ #define K32 (((w1s >> ( 0 + 7)) & 1) ? -1 : 0)
+ #define K33 (((w1s >> ( 0 + 6)) & 1) ? -1 : 0)
+ #define K34 (((w1s >> ( 0 + 5)) & 1) ? -1 : 0)
+ #define K35 (((w1s >> ( 0 + 4)) & 1) ? -1 : 0)
+ #define K36 (((w1s >> ( 0 + 3)) & 1) ? -1 : 0)
+ #define K37 (((w1s >> ( 0 + 2)) & 1) ? -1 : 0)
+ #define K38 (((w1s >> ( 0 + 1)) & 1) ? -1 : 0)
+ #define K39 (((w1s >> ( 0 + 0)) & 1) ? -1 : 0)
+ #define K40 (((w1s >> ( 8 + 7)) & 1) ? -1 : 0)
+ #define K41 (((w1s >> ( 8 + 6)) & 1) ? -1 : 0)
+ #define K42 (((w1s >> ( 8 + 5)) & 1) ? -1 : 0)
+ #define K43 (((w1s >> ( 8 + 4)) & 1) ? -1 : 0)
+ #define K44 (((w1s >> ( 8 + 3)) & 1) ? -1 : 0)
+ #define K45 (((w1s >> ( 8 + 2)) & 1) ? -1 : 0)
+ #define K46 (((w1s >> ( 8 + 1)) & 1) ? -1 : 0)
+ #define K47 (((w1s >> ( 8 + 0)) & 1) ? -1 : 0)
+ #define K48 (((w1s >> (16 + 7)) & 1) ? -1 : 0)
+ #define K49 (((w1s >> (16 + 6)) & 1) ? -1 : 0)
+ #define K50 (((w1s >> (16 + 5)) & 1) ? -1 : 0)
+ #define K51 (((w1s >> (16 + 4)) & 1) ? -1 : 0)
+ #define K52 (((w1s >> (16 + 3)) & 1) ? -1 : 0)
+ #define K53 (((w1s >> (16 + 2)) & 1) ? -1 : 0)
+ #define K54 (((w1s >> (16 + 1)) & 1) ? -1 : 0)
+ #define K55 (((w1s >> (16 + 0)) & 1) ? -1 : 0)
/**
* 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)
{
}
typedef VTYPE(ulong, VECT_SIZE) u64x;
#endif
-// this one needs to die
-#define allx(r) r
-
static inline u32 l32_from_64_S (u64 a)
{
const u32 r = (u32) (a);
static inline u64 rotr64_S (const u64 a, const u32 n)
{
- #if DEVICE_TYPE == DEVICE_TYPE_CPU
-
- const u64 r = rotate (a, (u64) 64 - n);
-
- #else
-
const u32 a0 = h32_from_64_S (a);
const u32 a1 = l32_from_64_S (a);
const u64 r = hl32_to_64_S (t0, t1);
- #endif
-
return r;
}
static inline u64x rotr64 (const u64x a, const u32 n)
{
- #if DEVICE_TYPE == DEVICE_TYPE_CPU
-
- const u64x r = rotate (a, (u64) 64 - n);
-
- #else
-
const u32x a0 = h32_from_64 (a);
const u32x a1 = l32_from_64 (a);
const u64x r = hl32_to_64 (t0, t1);
- #endif
-
return r;
}
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
* vendor specific
*/
-#if VENDOR_ID == 4098
+#if VENDOR_ID == (1 << 0)
#define IS_AMD
-#elif VENDOR_ID == 4318
+#elif VENDOR_ID == (1 << 6)
#define IS_NV
#else
#define IS_GENERIC
*/
#ifdef IS_AMD
+#ifndef cl_clang_storage_class_specifiers
#pragma OPENCL EXTENSION cl_amd_media_ops : enable
#pragma OPENCL EXTENSION cl_amd_media_ops2 : enable
#endif
+#endif
/**
* NV specific
* shared stuff
*/
-#define ETC_MAX (60 * 60 * 24 * 365 * 10)
+#define ETC_MAX (60 * 60 * 24 * 365 * 10)
-#define DEVICES_MAX 128
+#define DEVICES_MAX 128
-#define CL_PLATFORMS_MAX 16
+#define CL_PLATFORMS_MAX 16
-#define CL_VENDOR_NV "NVIDIA Corporation"
-#define CL_VENDOR_AMD "Advanced Micro Devices, Inc."
-#define CL_VENDOR_APPLE "Apple"
-#define CL_VENDOR_POCL "The pocl project"
+#define CL_VENDOR_AMD "Advanced Micro Devices, Inc."
+#define CL_VENDOR_APPLE "Apple"
+#define CL_VENDOR_INTEL_BEIGNET "Intel"
+#define CL_VENDOR_INTEL_SDK "Intel(R) OpenCL"
+#define CL_VENDOR_MESA "Mesa"
+#define CL_VENDOR_NV "NVIDIA Corporation"
+#define CL_VENDOR_POCL "The pocl project"
-#define VENDOR_ID_AMD 4098
-#define VENDOR_ID_NV 4318
-#define VENDOR_ID_APPLE_INTEL 4294967295
-#define VENDOR_ID_APPLE_IRIS 16925952
-#define VENDOR_ID_GENERIC 9999
+#define VENDOR_ID_AMD (1 << 0)
+#define VENDOR_ID_APPLE (1 << 1)
+#define VENDOR_ID_INTEL_BEIGNET (1 << 2)
+#define VENDOR_ID_INTEL_SDK (1 << 3)
+#define VENDOR_ID INTEL (1 << 4)
+#define VENDOR_ID_MESA (1 << 5)
+#define VENDOR_ID_NV (1 << 6)
+#define VENDOR_ID_POCL (1 << 7)
+#define VENDOR_ID_GENERIC (1 << 31)
-#define BLOCK_SIZE 64
+#define BLOCK_SIZE 64
-#define CHARSIZ 0x100
-#define INFOSZ CHARSIZ
+#define CHARSIZ 0x100
+#define INFOSZ CHARSIZ
-#define SP_HCSTAT "hashcat.hcstat"
-#define SP_PW_MIN 2
-#define SP_PW_MAX 64
-#define SP_ROOT_CNT (SP_PW_MAX * CHARSIZ)
-#define SP_MARKOV_CNT (SP_PW_MAX * CHARSIZ * CHARSIZ)
+#define SP_HCSTAT "hashcat.hcstat"
+#define SP_PW_MIN 2
+#define SP_PW_MAX 64
+#define SP_ROOT_CNT (SP_PW_MAX * CHARSIZ)
+#define SP_MARKOV_CNT (SP_PW_MAX * CHARSIZ * CHARSIZ)
-#define TUNING_DB_FILE "hashcat_tuning.hctab"
+#define TUNING_DB_FILE "hashcat_tuning.hctab"
-#define INDUCT_DIR "induct"
-#define OUTFILES_DIR "outfiles"
+#define INDUCT_DIR "induct"
+#define OUTFILES_DIR "outfiles"
-#define LOOPBACK_FILE "hashcat.loopback"
+#define LOOPBACK_FILE "hashcat.loopback"
-#define DICTSTAT_FILENAME "hashcat.dictstat"
-#define POTFILE_FILENAME "hashcat.pot"
+#define DICTSTAT_FILENAME "hashcat.dictstat"
+#define POTFILE_FILENAME "hashcat.pot"
/**
* types
*/
#ifdef _WIN
-typedef LARGE_INTEGER hc_timer_t;
-typedef HANDLE hc_thread_t;
-typedef CRITICAL_SECTION hc_thread_mutex_t;
+typedef LARGE_INTEGER hc_timer_t;
+typedef HANDLE hc_thread_t;
+typedef CRITICAL_SECTION hc_thread_mutex_t;
#elif _POSIX
-typedef struct timeval hc_timer_t;
-typedef pthread_t hc_thread_t;
-typedef pthread_mutex_t hc_thread_mutex_t;
+typedef struct timeval hc_timer_t;
+typedef pthread_t hc_thread_t;
+typedef pthread_mutex_t hc_thread_mutex_t;
#endif
#include <types.h>
#define DISPLAY_LEN_MAX_11100 10 + 32 + 1 + 8 + 1 + 32
#define DISPLAY_LEN_MIN_11200 9 + 40 + 1 + 40
#define DISPLAY_LEN_MAX_11200 9 + 40 + 1 + 40
-#define DISPLAY_LEN_MIN_11300 1 + 7 + 1 + 2 + 1 + 96 + 1 + 2 + 1 + 16 + 1 + 1 + 1 + 2 + 1 + 96 + 1 + 2 + 1 + 66
-#define DISPLAY_LEN_MAX_11300 1 + 7 + 1 + 2 + 1 + 96 + 1 + 2 + 1 + 16 + 1 + 6 + 1 + 2 + 1 + 96 + 1 + 2 + 1 + 66
+#define DISPLAY_LEN_MIN_11300 1 + 7 + 1 + 2 + 1 + 96 + 1 + 2 + 1 + 16 + 1 + 1 + 1 + 2 + 1 + 96 + 1 + 1 + 1 + 2
+#define DISPLAY_LEN_MAX_11300 1 + 7 + 1 + 2 + 1 + 96 + 1 + 2 + 1 + 16 + 1 + 6 + 1 + 2 + 1 + 96 + 1 + 3 + 1 + 512
#define DISPLAY_LEN_MIN_11400 6 + 0 + 1 + 0 + 1 + 0 + 1 + 0 + 1 + 0 + 1 + 0 + 1 + 1 + 1 + 0 + 1 + 1 + 1 + 0 + 1 + 0 + 1 + 0 + 1 + 3 + 1 + 32
#define DISPLAY_LEN_MAX_11400 6 + 512 + 1 + 512 + 1 + 116 + 1 + 116 + 1 + 246 + 1 + 245 + 1 + 246 + 1 + 245 + 1 + 50 + 1 + 50 + 1 + 50 + 1 + 50 + 1 + 3 + 1 + 32
#define DISPLAY_LEN_MIN_11500 8 + 1 + 8
uint exec_pos;
double exec_ms[EXEC_CACHE];
+ // this is "current" speed
+
uint speed_pos;
u64 speed_cnt[SPEED_CACHE];
- float speed_ms[SPEED_CACHE];
-
- hc_timer_t speed_rec[SPEED_CACHE];
+ double speed_ms[SPEED_CACHE];
hc_timer_t timer_speed;
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
#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;
-
- hc_timer_get (device_param->speed_rec[i], rec_ms);
-
- if (rec_ms > SPEED_MAXAGE) continue;
-
speed_cnt += device_param->speed_cnt[i];
speed_ms += device_param->speed_ms[i];
}
* 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++)
{
if (device_param->skipped) continue;
- // we need to clear values (set to 0) because in case the device does
- // not get new candidates it idles around but speed display would
- // show it as working.
- // if we instantly set it to 0 after reading it happens that the
- // speed can be shown as zero if the users refreshes too fast.
- // therefore, we add a timestamp when a stat was recorded and if its
- // too old we will not use it
-
speed_cnt[device_id] = 0;
speed_ms[device_id] = 0;
for (int i = 0; i < SPEED_CACHE; i++)
{
- float rec_ms;
-
- hc_timer_get (device_param->speed_rec[i], rec_ms);
-
- if (rec_ms > SPEED_MAXAGE) continue;
-
speed_cnt[device_id] += device_param->speed_cnt[i];
speed_ms[device_id] += device_param->speed_ms[i];
}
* 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.devices_status == STATUS_CRACKED) break;
if (data.devices_status == STATUS_ABORTED) break;
if (data.devices_status == STATUS_QUIT) break;
+
+ /**
+ * speed
+ */
+
+ const float iter_part = (float) (loop_pos + loop_left) / iter;
+
+ const u64 perf_sum_all = pws_cnt * iter_part;
+
+ double speed_ms;
+
+ hc_timer_get (device_param->timer_speed, speed_ms);
+
+ const u32 speed_pos = device_param->speed_pos;
+
+ device_param->speed_cnt[speed_pos] = perf_sum_all;
+
+ device_param->speed_ms[speed_pos] = speed_ms;
}
if (opts_type & OPTS_TYPE_HOOK23)
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;
- device_param->speed_rec[speed_pos] = device_param->timer_speed;
-
hc_thread_mutex_unlock (mux_display);
speed_pos++;
switch (hash_mode)
{
- case 1500: hashes_buf[0].salt->salt_len = 2;
+ case 1500: hashes_buf[0].salt->salt_len = 2;
+ hashes_buf[0].salt->salt_buf[0] = 388; // pure magic
break;
case 1731: hashes_buf[0].salt->salt_len = 4;
break;
hc_clGetDeviceIDs (data.ocl, platform, CL_DEVICE_TYPE_ALL, DEVICES_MAX, platform_devices, &platform_devices_cnt);
+ char platform_vendor[INFOSZ] = { 0 };
+
+ hc_clGetPlatformInfo (data.ocl, platform, CL_PLATFORM_VENDOR, sizeof (platform_vendor), platform_vendor, NULL);
+
+ // find our own platform vendor because pocl and mesa are pushing original vendor_id through opencl
+ // this causes trouble with vendor id based macros
+ // we'll assign generic to those without special optimization available
+
+ cl_uint vendor_id = 0;
+
+ if (strcmp (platform_vendor, CL_VENDOR_AMD) == 0)
+ {
+ vendor_id = VENDOR_ID_AMD;
+ }
+ else if (strcmp (platform_vendor, CL_VENDOR_APPLE) == 0)
+ {
+ vendor_id = VENDOR_ID_GENERIC;
+ }
+ else if (strcmp (platform_vendor, CL_VENDOR_INTEL_BEIGNET) == 0)
+ {
+ vendor_id = VENDOR_ID_GENERIC;
+ }
+ else if (strcmp (platform_vendor, CL_VENDOR_INTEL_SDK) == 0)
+ {
+ vendor_id = VENDOR_ID_GENERIC;
+ }
+ else if (strcmp (platform_vendor, CL_VENDOR_MESA) == 0)
+ {
+ vendor_id = VENDOR_ID_GENERIC;
+ }
+ else if (strcmp (platform_vendor, CL_VENDOR_NV) == 0)
+ {
+ vendor_id = VENDOR_ID_NV;
+ }
+ else if (strcmp (platform_vendor, CL_VENDOR_POCL) == 0)
+ {
+ vendor_id = VENDOR_ID_GENERIC;
+ }
+ else
+ {
+ vendor_id = VENDOR_ID_GENERIC;
+ }
+
for (uint platform_devices_id = 0; platform_devices_id < platform_devices_cnt; platform_devices_id++)
{
size_t param_value_size = 0;
hc_device_param_t *device_param = &data.devices_param[device_id];
+ device_param->vendor_id = vendor_id;
+
device_param->device = platform_devices[platform_devices_id];
device_param->device_id = device_id;
device_param->device_type = device_type;
- // vendor_id
-
- cl_uint vendor_id = 0;
-
- hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VENDOR_ID, sizeof (vendor_id), &vendor_id, NULL);
-
- device_param->vendor_id = vendor_id;
-
// device_name
hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_NAME, 0, NULL, ¶m_value_size);
myfree (device_opencl_version);
- if (strstr (device_version, "pocl"))
- {
- // pocl returns the real vendor_id in CL_DEVICE_VENDOR_ID which causes many problems because of hms and missing amd_bfe () etc
- // we need to overwrite vendor_id to avoid this. maybe open pocl issue?
-
- cl_uint vendor_id = VENDOR_ID_GENERIC;
-
- device_param->vendor_id = vendor_id;
- }
-
// vector_width
cl_uint vector_width;
if (device_param->skipped == 0)
{
- if (strstr (device_version, "pocl"))
- {
- if (force == 0)
- {
- log_info ("");
- log_info ("ATTENTION! All pocl drivers are known to be broken due to broken LLVM <= 3.7");
- log_info ("You are STRONGLY encouraged not to use it");
- log_info ("You can use --force to override this but do not post error reports if you do so");
- log_info ("");
-
- return (-1);
- }
- }
-
if (device_type & CL_DEVICE_TYPE_GPU)
{
- if (vendor_id == VENDOR_ID_NV)
- {
- if (device_param->kernel_exec_timeout != 0)
- {
- if (data.quiet == 0) log_info ("Device #%u: WARNING! Kernel exec timeout is not disabled, it might cause you errors of code 702", device_id + 1);
- if (data.quiet == 0) log_info (" See the wiki on how to disable it: https://hashcat.net/wiki/doku.php?id=timeout_patch");
- }
- }
- else if (vendor_id == VENDOR_ID_AMD)
+ if (vendor_id == VENDOR_ID_AMD)
{
int catalyst_check = (force == 1) ? 0 : 1;
log_info ("You can use --force to override this but do not post error reports if you do so");
log_info ("");
+ return (-1);
+ }
+ }
+ else if (vendor_id == VENDOR_ID_NV)
+ {
+ if (device_param->kernel_exec_timeout != 0)
+ {
+ if (data.quiet == 0) log_info ("Device #%u: WARNING! Kernel exec timeout is not disabled, it might cause you errors of code 702", device_id + 1);
+ if (data.quiet == 0) log_info (" See the wiki on how to disable it: https://hashcat.net/wiki/doku.php?id=timeout_patch");
+ }
+ }
+ else if (vendor_id == VENDOR_ID_POCL)
+ {
+ if (force == 0)
+ {
+ log_info ("");
+ log_info ("ATTENTION! All pocl drivers are known to be broken due to broken LLVM <= 3.7");
+ log_info ("You are STRONGLY encouraged not to use it");
+ log_info ("You can use --force to override this but do not post error reports if you do so");
+ log_info ("");
+
return (-1);
}
}
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_rec, 0, SPEED_CACHE * sizeof (hc_timer_t));
+ memset (device_param->speed_ms, 0, SPEED_CACHE * sizeof (double));
device_param->exec_pos = 0;
{
if (data.devices_status == STATUS_PAUSED)
{
- float ms_paused;
+ double ms_paused;
hc_timer_get (data.timer_paused, ms_paused);
projects / hashcat.git / commitdiff