Jens Steube 7 years ago
parent 922fea7616
commit 4a3c90dd3c

@ -76,10 +76,20 @@ inline u32 check (const u32 digest[4], __global const u32 *bitmap_s1_a, __global
return (1);
}
inline void mark_hash (__global plain_t *plains_buf, __global u32 *d_result, const int salt_pos, const int digest_pos, const int hash_pos, const u32 gid, const u32 il_pos)
inline void mark_hash (__global plain_t *plains_buf, __global u32 *d_result, const u32 salt_pos, const u32 digests_cnt, const u32 digest_pos, const u32 hash_pos, const u32 gid, const u32 il_pos)
{
const u32 idx = atomic_inc (d_result);
if (idx >= digests_cnt)
{
// this is kind of tricky: we *must* call atomic_inc() to know about the current value from a multi-thread perspective
// this action creates a buffer overflow, so we need to fix it here
atomic_dec (d_result);
return;
}
plains_buf[idx].salt_pos = salt_pos;
plains_buf[idx].digest_pos = digest_pos; // relative
plains_buf[idx].hash_pos = hash_pos; // absolute

@ -26,7 +26,7 @@ if (check (digest_tp,
if (atomic_inc (&hashes_shown[final_hash_pos]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos);
}
}
}

@ -28,7 +28,7 @@ if (check (digest_tp,
if (atomic_inc (&hashes_shown[final_hash_pos]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + slice);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + slice);
}
}
}

@ -7,6 +7,6 @@ if ((r0 == search[0])
if (atomic_inc (&hashes_shown[final_hash_pos]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos);
}
}

@ -5,6 +5,6 @@ if ((il_pos + slice) < il_cnt)
if (atomic_inc (&hashes_shown[final_hash_pos]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + slice);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + slice);
}
}

@ -18,7 +18,7 @@
\
if (atomic_inc (&hashes_shown[final_hash_pos]) == 0) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos); \
} \
} \
}
@ -42,7 +42,7 @@
\
if (atomic_inc (&hashes_shown[final_hash_pos]) == 0) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos); \
} \
} \
} \
@ -67,7 +67,7 @@
\
if (vector_accessible (il_pos, il_cnt, 0) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 0); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 0); \
} \
} \
\
@ -77,7 +77,7 @@
\
if (vector_accessible (il_pos, il_cnt, 1) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 1); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 1); \
} \
} \
}
@ -102,7 +102,7 @@
\
if (vector_accessible (il_pos, il_cnt, 0) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 0); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 0); \
} \
} \
} \
@ -122,7 +122,7 @@
\
if (vector_accessible (il_pos, il_cnt, 1) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 1); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 1); \
} \
} \
} \
@ -145,7 +145,7 @@
\
if (vector_accessible (il_pos, il_cnt, 0) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 0); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 0); \
} \
} \
\
@ -155,7 +155,7 @@
\
if (vector_accessible (il_pos, il_cnt, 1) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 1); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 1); \
} \
} \
\
@ -165,7 +165,7 @@
\
if (vector_accessible (il_pos, il_cnt, 2) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 2); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 2); \
} \
} \
\
@ -175,7 +175,7 @@
\
if (vector_accessible (il_pos, il_cnt, 3) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 3); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 3); \
} \
} \
}
@ -202,7 +202,7 @@
\
if (vector_accessible (il_pos, il_cnt, 0) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 0); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 0); \
} \
} \
} \
@ -222,7 +222,7 @@
\
if (vector_accessible (il_pos, il_cnt, 1) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 1); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 1); \
} \
} \
} \
@ -242,7 +242,7 @@
\
if (vector_accessible (il_pos, il_cnt, 2) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 2); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 2); \
} \
} \
} \
@ -262,7 +262,7 @@
\
if (vector_accessible (il_pos, il_cnt, 3) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 3); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 3); \
} \
} \
} \
@ -285,7 +285,7 @@
\
if (vector_accessible (il_pos, il_cnt, 0) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 0); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 0); \
} \
} \
\
@ -295,7 +295,7 @@
\
if (vector_accessible (il_pos, il_cnt, 1) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 1); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 1); \
} \
} \
\
@ -305,7 +305,7 @@
\
if (vector_accessible (il_pos, il_cnt, 2) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 2); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 2); \
} \
} \
\
@ -315,7 +315,7 @@
\
if (vector_accessible (il_pos, il_cnt, 3) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 3); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 3); \
} \
} \
if (((h0).s4 == search[0]) && ((h1).s4 == search[1]) && ((h2).s4 == search[2]) && ((h3).s4 == search[3])) \
@ -324,7 +324,7 @@
\
if (vector_accessible (il_pos, il_cnt, 4) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 4); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 4); \
} \
} \
\
@ -334,7 +334,7 @@
\
if (vector_accessible (il_pos, il_cnt, 5) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 5); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 5); \
} \
} \
\
@ -344,7 +344,7 @@
\
if (vector_accessible (il_pos, il_cnt, 6) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 6); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 6); \
} \
} \
\
@ -354,7 +354,7 @@
\
if (vector_accessible (il_pos, il_cnt, 7) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 7); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 7); \
} \
} \
}
@ -385,7 +385,7 @@
\
if (vector_accessible (il_pos, il_cnt, 0) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 0); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 0); \
} \
} \
} \
@ -405,7 +405,7 @@
\
if (vector_accessible (il_pos, il_cnt, 1) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 1); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 1); \
} \
} \
} \
@ -425,7 +425,7 @@
\
if (vector_accessible (il_pos, il_cnt, 2) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 2); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 2); \
} \
} \
} \
@ -445,7 +445,7 @@
\
if (vector_accessible (il_pos, il_cnt, 3) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 3); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 3); \
} \
} \
} \
@ -464,7 +464,7 @@
\
if (vector_accessible (il_pos, il_cnt, 4) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 4); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 4); \
} \
} \
} \
@ -484,7 +484,7 @@
\
if (vector_accessible (il_pos, il_cnt, 5) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 5); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 5); \
} \
} \
} \
@ -504,7 +504,7 @@
\
if (vector_accessible (il_pos, il_cnt, 6) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 6); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 6); \
} \
} \
} \
@ -524,7 +524,7 @@
\
if (vector_accessible (il_pos, il_cnt, 7) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 7); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 7); \
} \
} \
} \
@ -547,7 +547,7 @@
\
if (vector_accessible (il_pos, il_cnt, 0) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 0); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 0); \
} \
} \
\
@ -557,7 +557,7 @@
\
if (vector_accessible (il_pos, il_cnt, 1) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 1); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 1); \
} \
} \
\
@ -567,7 +567,7 @@
\
if (vector_accessible (il_pos, il_cnt, 2) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 2); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 2); \
} \
} \
\
@ -577,7 +577,7 @@
\
if (vector_accessible (il_pos, il_cnt, 3) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 3); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 3); \
} \
} \
if (((h0).s4 == search[0]) && ((h1).s4 == search[1]) && ((h2).s4 == search[2]) && ((h3).s4 == search[3])) \
@ -586,7 +586,7 @@
\
if (vector_accessible (il_pos, il_cnt, 4) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 4); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 4); \
} \
} \
\
@ -596,7 +596,7 @@
\
if (vector_accessible (il_pos, il_cnt, 5) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 5); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 5); \
} \
} \
\
@ -606,7 +606,7 @@
\
if (vector_accessible (il_pos, il_cnt, 6) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 6); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 6); \
} \
} \
\
@ -616,7 +616,7 @@
\
if (vector_accessible (il_pos, il_cnt, 7) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 7); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 7); \
} \
} \
\
@ -626,7 +626,7 @@
\
if (vector_accessible (il_pos, il_cnt, 8) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 8); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 8); \
} \
} \
\
@ -636,7 +636,7 @@
\
if (vector_accessible (il_pos, il_cnt, 9) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 9); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 9); \
} \
} \
\
@ -646,7 +646,7 @@
\
if (vector_accessible (il_pos, il_cnt, 10) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 10); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 10); \
} \
} \
\
@ -656,7 +656,7 @@
\
if (vector_accessible (il_pos, il_cnt, 11) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 11); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 11); \
} \
} \
\
@ -666,7 +666,7 @@
\
if (vector_accessible (il_pos, il_cnt, 12) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 12); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 12); \
} \
} \
\
@ -676,7 +676,7 @@
\
if (vector_accessible (il_pos, il_cnt, 13) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 13); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 13); \
} \
} \
\
@ -686,7 +686,7 @@
\
if (vector_accessible (il_pos, il_cnt, 14) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 14); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 14); \
} \
} \
\
@ -696,7 +696,7 @@
\
if (vector_accessible (il_pos, il_cnt, 15) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + 15); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos + 15); \
} \
} \
}
@ -735,7 +735,7 @@
\
if (vector_accessible (il_pos, il_cnt, 0) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 0); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 0); \
} \
} \
} \
@ -755,7 +755,7 @@
\
if (vector_accessible (il_pos, il_cnt, 1) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 1); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 1); \
} \
} \
} \
@ -775,7 +775,7 @@
\
if (vector_accessible (il_pos, il_cnt, 2) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 2); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 2); \
} \
} \
} \
@ -795,7 +795,7 @@
\
if (vector_accessible (il_pos, il_cnt, 3) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 3); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 3); \
} \
} \
} \
@ -815,7 +815,7 @@
\
if (vector_accessible (il_pos, il_cnt, 4) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 4); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 4); \
} \
} \
} \
@ -835,7 +835,7 @@
\
if (vector_accessible (il_pos, il_cnt, 5) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 5); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 5); \
} \
} \
} \
@ -855,7 +855,7 @@
\
if (vector_accessible (il_pos, il_cnt, 6) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 6); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 6); \
} \
} \
} \
@ -875,7 +875,7 @@
\
if (vector_accessible (il_pos, il_cnt, 7) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 7); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 7); \
} \
} \
} \
@ -895,7 +895,7 @@
\
if (vector_accessible (il_pos, il_cnt, 8) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 8); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 8); \
} \
} \
} \
@ -915,7 +915,7 @@
\
if (vector_accessible (il_pos, il_cnt, 9) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 9); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 9); \
} \
} \
} \
@ -935,7 +935,7 @@
\
if (vector_accessible (il_pos, il_cnt, 10) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 10); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 10); \
} \
} \
} \
@ -955,7 +955,7 @@
\
if (vector_accessible (il_pos, il_cnt, 11) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 11); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 11); \
} \
} \
} \
@ -975,7 +975,7 @@
\
if (vector_accessible (il_pos, il_cnt, 12) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 12); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 12); \
} \
} \
} \
@ -995,7 +995,7 @@
\
if (vector_accessible (il_pos, il_cnt, 13) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 13); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 13); \
} \
} \
} \
@ -1015,7 +1015,7 @@
\
if (vector_accessible (il_pos, il_cnt, 14) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 14); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 14); \
} \
} \
} \
@ -1035,7 +1035,7 @@
\
if (vector_accessible (il_pos, il_cnt, 15) && (atomic_inc (&hashes_shown[final_hash_pos]) == 0)) \
{ \
mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + 15); \
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos + 15); \
} \
} \
} \

@ -692,16 +692,16 @@ __kernel void m06211_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -692,17 +692,17 @@ __kernel void m06212_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
u32 ukey3[8];
@ -729,16 +729,16 @@ __kernel void m06212_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -692,17 +692,17 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
u32 ukey3[8];
@ -733,17 +733,17 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
#if defined (IS_APPLE) && defined (IS_GPU)
@ -778,11 +778,11 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent_twofish_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -600,16 +600,16 @@ __kernel void m06221_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -600,17 +600,17 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
u32 ukey3[8];
@ -637,16 +637,16 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -649,17 +649,17 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
u32 ukey3[8];
@ -690,17 +690,17 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
volatile u32 ukey5[8];
@ -727,11 +727,11 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent_twofish_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -2235,16 +2235,16 @@ __kernel void m06231_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -2004,17 +2004,17 @@ __kernel void m06232_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
u32 ukey3[8];
@ -2041,16 +2041,16 @@ __kernel void m06232_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -2004,17 +2004,17 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
u32 ukey3[8];
@ -2045,17 +2045,17 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
#if defined (IS_APPLE) && defined (IS_GPU)
@ -2090,11 +2090,11 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent_twofish_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -1565,7 +1565,7 @@ __kernel void m06800_comp (__global pw_t *pws, __global const kernel_rule_t *rul
&& (out[2] == salt_buf[2])
&& (out[3] == salt_buf[3]))
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0);
}
/**

@ -634,7 +634,7 @@ __kernel void m07500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}
}
@ -735,7 +735,7 @@ __kernel void m07500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}
}

@ -682,7 +682,7 @@ __kernel void m07500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}
}
@ -833,7 +833,7 @@ __kernel void m07500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}
}

@ -626,7 +626,7 @@ void m07500 (__local RC4_KEY *rc4_keys, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[
if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}
}

@ -1885,7 +1885,7 @@ __kernel void m08800_comp (__global pw_t *pws, __global const kernel_rule_t *rul
// MSDOS5.0
if ((r0 == 0x4f44534d) && (r1 == 0x302e3553))
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0);
}
}
@ -1951,7 +1951,7 @@ __kernel void m08800_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if ((r[5] < 2) && (r[6] < 16) && ((r[14] & 0xffff) == 0xEF53))
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0);
}
}
}

@ -1341,6 +1341,6 @@ __kernel void m11300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
&& (out[2] == 0x10101010)
&& (out[3] == 0x10101010))
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0);
}
}

@ -853,7 +853,7 @@ __kernel void m11600_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (seven_zip_hook[gid].hook_success == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0);
return;
}

@ -770,7 +770,7 @@ void m13100 (__local RC4_KEY *rc4_keys, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[
if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[salt_pos].edata2, krb5tgs_bufs[salt_pos].edata2_len, K2, checksum) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}
}
@ -855,7 +855,7 @@ __kernel void m13100_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[salt_pos].edata2, krb5tgs_bufs[salt_pos].edata2_len, K2, checksum) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}
}
@ -948,7 +948,7 @@ __kernel void m13100_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[salt_pos].edata2, krb5tgs_bufs[salt_pos].edata2_len, K2, checksum) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}
}

@ -850,7 +850,7 @@ __kernel void m13100_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[salt_pos].edata2, krb5tgs_bufs[salt_pos].edata2_len, K2, checksum) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}
}
@ -992,7 +992,7 @@ __kernel void m13100_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[salt_pos].edata2, krb5tgs_bufs[salt_pos].edata2_len, K2, checksum) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}
}

@ -772,7 +772,7 @@ void m13100 (__local RC4_KEY *rc4_keys, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[
if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[salt_pos].edata2, krb5tgs_bufs[salt_pos].edata2_len, K2, checksum) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}
}

@ -1211,6 +1211,6 @@ __kernel void m13200_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if(tmps[gid].cipher[0]==0xA6A6A6A6 && tmps[gid].cipher[1]==0xA6A6A6A6)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}

@ -994,7 +994,7 @@ __kernel void m13400_comp (__global pw_t *pws, __global const kernel_rule_t *rul
&& esalt_bufs[salt_pos].contents_hash[6] == final_digest[6]
&& esalt_bufs[salt_pos].contents_hash[7] == final_digest[7])
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}
else
@ -1198,7 +1198,7 @@ __kernel void m13400_comp (__global pw_t *pws, __global const kernel_rule_t *rul
&& esalt_bufs[salt_pos].contents_hash[6] == final_digest[6]
&& esalt_bufs[salt_pos].contents_hash[7] == final_digest[7])
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}
}
@ -1236,7 +1236,7 @@ __kernel void m13400_comp (__global pw_t *pws, __global const kernel_rule_t *rul
&& esalt_bufs[salt_pos].expected_bytes[2] == out[2]
&& esalt_bufs[salt_pos].expected_bytes[3] == out[3])
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}
}

@ -663,16 +663,16 @@ __kernel void m13751_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -663,17 +663,17 @@ __kernel void m13752_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
u32 ukey3[8];
@ -700,16 +700,16 @@ __kernel void m13752_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -663,17 +663,17 @@ __kernel void m13753_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
u32 ukey3[8];
@ -700,17 +700,17 @@ __kernel void m13753_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
u32 ukey5[8];
@ -737,11 +737,11 @@ __kernel void m13753_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (verify_header_aes_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
if (verify_header_serpent_twofish_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -701,6 +701,6 @@ __kernel void m14611_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (entropy < MAX_ENTROPY)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -648,6 +648,6 @@ __kernel void m14612_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (entropy < MAX_ENTROPY)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -648,6 +648,6 @@ __kernel void m14613_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (entropy < MAX_ENTROPY)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -692,6 +692,6 @@ __kernel void m14621_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (entropy < MAX_ENTROPY)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -639,6 +639,6 @@ __kernel void m14622_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (entropy < MAX_ENTROPY)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -639,6 +639,6 @@ __kernel void m14623_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (entropy < MAX_ENTROPY)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -668,6 +668,6 @@ __kernel void m14631_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (entropy < MAX_ENTROPY)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -615,6 +615,6 @@ __kernel void m14632_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (entropy < MAX_ENTROPY)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -615,6 +615,6 @@ __kernel void m14633_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (entropy < MAX_ENTROPY)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -856,6 +856,6 @@ __kernel void m14641_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (entropy < MAX_ENTROPY)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -803,6 +803,6 @@ __kernel void m14642_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (entropy < MAX_ENTROPY)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -803,6 +803,6 @@ __kernel void m14643_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (entropy < MAX_ENTROPY)
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}

@ -1742,7 +1742,7 @@ __kernel void m14700_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if ((cipher[0] == 0xa6a6a6a6) && (cipher[1] == 0xa6a6a6a6))
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0);
return;
}

@ -2297,7 +2297,7 @@ __kernel void m14800_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if ((cipher[0] == 0xa6a6a6a6) && (cipher[1] == 0xa6a6a6a6))
{
mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, 0);
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0);
return;
}

Loading…
Cancel
Save