diff --git a/OpenCL/inc_common.cl b/OpenCL/inc_common.cl index abe27347a..de269e471 100644 --- a/OpenCL/inc_common.cl +++ b/OpenCL/inc_common.cl @@ -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 diff --git a/OpenCL/inc_comp_multi.cl b/OpenCL/inc_comp_multi.cl index 14ad55cc0..b10bfce00 100644 --- a/OpenCL/inc_comp_multi.cl +++ b/OpenCL/inc_comp_multi.cl @@ -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); } } } diff --git a/OpenCL/inc_comp_multi_bs.cl b/OpenCL/inc_comp_multi_bs.cl index 894416a17..8534a0467 100644 --- a/OpenCL/inc_comp_multi_bs.cl +++ b/OpenCL/inc_comp_multi_bs.cl @@ -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); } } } diff --git a/OpenCL/inc_comp_single.cl b/OpenCL/inc_comp_single.cl index 631a8bc17..7895d50cc 100644 --- a/OpenCL/inc_comp_single.cl +++ b/OpenCL/inc_comp_single.cl @@ -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); } } diff --git a/OpenCL/inc_comp_single_bs.cl b/OpenCL/inc_comp_single_bs.cl index 92383d6e1..02ff059eb 100644 --- a/OpenCL/inc_comp_single_bs.cl +++ b/OpenCL/inc_comp_single_bs.cl @@ -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); } } diff --git a/OpenCL/inc_simd.cl b/OpenCL/inc_simd.cl index 2e8b3cdee..fc62b9c96 100644 --- a/OpenCL/inc_simd.cl +++ b/OpenCL/inc_simd.cl @@ -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); \ } \ } \ } \ diff --git a/OpenCL/m06211.cl b/OpenCL/m06211.cl index d69d5bd2e..aeaa6c76b 100644 --- a/OpenCL/m06211.cl +++ b/OpenCL/m06211.cl @@ -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); } } diff --git a/OpenCL/m06212.cl b/OpenCL/m06212.cl index e9d6b2a4a..ad797a439 100644 --- a/OpenCL/m06212.cl +++ b/OpenCL/m06212.cl @@ -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); } } diff --git a/OpenCL/m06213.cl b/OpenCL/m06213.cl index 6f40528a8..6b345796b 100644 --- a/OpenCL/m06213.cl +++ b/OpenCL/m06213.cl @@ -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); } } diff --git a/OpenCL/m06221.cl b/OpenCL/m06221.cl index de43c202e..dedca8b34 100644 --- a/OpenCL/m06221.cl +++ b/OpenCL/m06221.cl @@ -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); } } diff --git a/OpenCL/m06222.cl b/OpenCL/m06222.cl index b931d6d1f..797bd191b 100644 --- a/OpenCL/m06222.cl +++ b/OpenCL/m06222.cl @@ -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); } } diff --git a/OpenCL/m06223.cl b/OpenCL/m06223.cl index 096ec23d2..02adac994 100644 --- a/OpenCL/m06223.cl +++ b/OpenCL/m06223.cl @@ -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); } } diff --git a/OpenCL/m06231.cl b/OpenCL/m06231.cl index 35b5219cc..a5530aa8a 100644 --- a/OpenCL/m06231.cl +++ b/OpenCL/m06231.cl @@ -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); } } diff --git a/OpenCL/m06232.cl b/OpenCL/m06232.cl index 1735fd340..38a6acc6e 100644 --- a/OpenCL/m06232.cl +++ b/OpenCL/m06232.cl @@ -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); } } diff --git a/OpenCL/m06233.cl b/OpenCL/m06233.cl index 75f565cdc..666d35f86 100644 --- a/OpenCL/m06233.cl +++ b/OpenCL/m06233.cl @@ -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); } } diff --git a/OpenCL/m06800.cl b/OpenCL/m06800.cl index 54253014b..951ce06c4 100644 --- a/OpenCL/m06800.cl +++ b/OpenCL/m06800.cl @@ -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); } /** diff --git a/OpenCL/m07500_a0.cl b/OpenCL/m07500_a0.cl index f62e80127..5f4c6da38 100644 --- a/OpenCL/m07500_a0.cl +++ b/OpenCL/m07500_a0.cl @@ -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); } } } diff --git a/OpenCL/m07500_a1.cl b/OpenCL/m07500_a1.cl index efdf582f5..e78cf7a6c 100644 --- a/OpenCL/m07500_a1.cl +++ b/OpenCL/m07500_a1.cl @@ -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); } } } diff --git a/OpenCL/m07500_a3.cl b/OpenCL/m07500_a3.cl index 9a05d4837..73fe24bac 100644 --- a/OpenCL/m07500_a3.cl +++ b/OpenCL/m07500_a3.cl @@ -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); } } } diff --git a/OpenCL/m08800.cl b/OpenCL/m08800.cl index 79e54f1f1..9a0a491e9 100644 --- a/OpenCL/m08800.cl +++ b/OpenCL/m08800.cl @@ -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); } } } diff --git a/OpenCL/m11300.cl b/OpenCL/m11300.cl index f85c79b0a..f42dd66cd 100644 --- a/OpenCL/m11300.cl +++ b/OpenCL/m11300.cl @@ -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); } } diff --git a/OpenCL/m11600.cl b/OpenCL/m11600.cl index 581d44bae..00006d356 100644 --- a/OpenCL/m11600.cl +++ b/OpenCL/m11600.cl @@ -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; } diff --git a/OpenCL/m13100_a0.cl b/OpenCL/m13100_a0.cl index b173ba127..cf60fd0a5 100644 --- a/OpenCL/m13100_a0.cl +++ b/OpenCL/m13100_a0.cl @@ -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); } } } diff --git a/OpenCL/m13100_a1.cl b/OpenCL/m13100_a1.cl index 571803ea8..91e237ad8 100644 --- a/OpenCL/m13100_a1.cl +++ b/OpenCL/m13100_a1.cl @@ -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); } } } diff --git a/OpenCL/m13100_a3.cl b/OpenCL/m13100_a3.cl index 18610708d..fbaa9c4bf 100644 --- a/OpenCL/m13100_a3.cl +++ b/OpenCL/m13100_a3.cl @@ -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); } } } diff --git a/OpenCL/m13200.cl b/OpenCL/m13200.cl index 47c36826b..586c08b65 100644 --- a/OpenCL/m13200.cl +++ b/OpenCL/m13200.cl @@ -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); } } diff --git a/OpenCL/m13400.cl b/OpenCL/m13400.cl index 6993a40c6..fb1700188 100644 --- a/OpenCL/m13400.cl +++ b/OpenCL/m13400.cl @@ -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); } } } diff --git a/OpenCL/m13751.cl b/OpenCL/m13751.cl index efcd829db..5e52fec2f 100644 --- a/OpenCL/m13751.cl +++ b/OpenCL/m13751.cl @@ -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); } } diff --git a/OpenCL/m13752.cl b/OpenCL/m13752.cl index 4aca26385..b0d1a87b8 100644 --- a/OpenCL/m13752.cl +++ b/OpenCL/m13752.cl @@ -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); } } diff --git a/OpenCL/m13753.cl b/OpenCL/m13753.cl index 0784e1d35..f312d1dfa 100644 --- a/OpenCL/m13753.cl +++ b/OpenCL/m13753.cl @@ -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); } } diff --git a/OpenCL/m14611.cl b/OpenCL/m14611.cl index 86c70b2a2..56442970b 100644 --- a/OpenCL/m14611.cl +++ b/OpenCL/m14611.cl @@ -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); } } diff --git a/OpenCL/m14612.cl b/OpenCL/m14612.cl index 0b7fcb8d0..bbc616584 100644 --- a/OpenCL/m14612.cl +++ b/OpenCL/m14612.cl @@ -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); } } diff --git a/OpenCL/m14613.cl b/OpenCL/m14613.cl index d1ca62d38..d065b907a 100644 --- a/OpenCL/m14613.cl +++ b/OpenCL/m14613.cl @@ -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); } } diff --git a/OpenCL/m14621.cl b/OpenCL/m14621.cl index 991812d89..630567b52 100644 --- a/OpenCL/m14621.cl +++ b/OpenCL/m14621.cl @@ -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); } } diff --git a/OpenCL/m14622.cl b/OpenCL/m14622.cl index 3fa60e145..250f7142c 100644 --- a/OpenCL/m14622.cl +++ b/OpenCL/m14622.cl @@ -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); } } diff --git a/OpenCL/m14623.cl b/OpenCL/m14623.cl index 739416602..1aa93e2ff 100644 --- a/OpenCL/m14623.cl +++ b/OpenCL/m14623.cl @@ -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); } } diff --git a/OpenCL/m14631.cl b/OpenCL/m14631.cl index 95d395319..812f5052c 100644 --- a/OpenCL/m14631.cl +++ b/OpenCL/m14631.cl @@ -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); } } diff --git a/OpenCL/m14632.cl b/OpenCL/m14632.cl index 05f167077..d8e4202f9 100644 --- a/OpenCL/m14632.cl +++ b/OpenCL/m14632.cl @@ -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); } } diff --git a/OpenCL/m14633.cl b/OpenCL/m14633.cl index 88ef85648..7f3dfe0b6 100644 --- a/OpenCL/m14633.cl +++ b/OpenCL/m14633.cl @@ -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); } } diff --git a/OpenCL/m14641.cl b/OpenCL/m14641.cl index 751acfc72..6fa65e82f 100644 --- a/OpenCL/m14641.cl +++ b/OpenCL/m14641.cl @@ -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); } } diff --git a/OpenCL/m14642.cl b/OpenCL/m14642.cl index 7b21f97a4..5de0e1c0f 100644 --- a/OpenCL/m14642.cl +++ b/OpenCL/m14642.cl @@ -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); } } diff --git a/OpenCL/m14643.cl b/OpenCL/m14643.cl index b68eea37d..002ec4a3e 100644 --- a/OpenCL/m14643.cl +++ b/OpenCL/m14643.cl @@ -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); } } diff --git a/OpenCL/m14700.cl b/OpenCL/m14700.cl index cc69ee987..b802c5cf8 100644 --- a/OpenCL/m14700.cl +++ b/OpenCL/m14700.cl @@ -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; } diff --git a/OpenCL/m14800.cl b/OpenCL/m14800.cl index e21ecc83a..f52fa0dbe 100644 --- a/OpenCL/m14800.cl +++ b/OpenCL/m14800.cl @@ -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; }