diff --git a/OpenCL/check_multi_comp4.c b/OpenCL/check_multi_comp4.c index 6b384162b..f45a03b41 100644 --- a/OpenCL/check_multi_comp4.c +++ b/OpenCL/check_multi_comp4.c @@ -26,7 +26,7 @@ if (check (digest_tp, if (atomic_add (&hashes_shown[final_hash_pos], 1) == 0) { - mark_hash (plains_buf, d_result, salt_pos, digest_pos, final_hash_pos, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos); } } } diff --git a/OpenCL/check_multi_comp4_bs.c b/OpenCL/check_multi_comp4_bs.c index a2b371d5e..0fbcd7bf3 100644 --- a/OpenCL/check_multi_comp4_bs.c +++ b/OpenCL/check_multi_comp4_bs.c @@ -26,7 +26,7 @@ if (check (digest_tp, if (atomic_add (&hashes_shown[final_hash_pos], 1) == 0) { - mark_hash (plains_buf, d_result, salt_pos, digest_pos, final_hash_pos, gid, il_pos + slice); + mark_hash (plains_buf, d_return_buf, salt_pos, digest_pos, final_hash_pos, gid, il_pos + slice); } } } diff --git a/OpenCL/check_single_comp4.c b/OpenCL/check_single_comp4.c index 7eba5e201..f754d8e9b 100644 --- a/OpenCL/check_single_comp4.c +++ b/OpenCL/check_single_comp4.c @@ -7,6 +7,6 @@ if ((r0 == search[0]) if (atomic_add (&hashes_shown[final_hash_pos], 1) == 0) { - mark_hash (plains_buf, d_result, salt_pos, 0, final_hash_pos, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos); } } diff --git a/OpenCL/check_single_comp4_bs.c b/OpenCL/check_single_comp4_bs.c index f2c5366fc..6ec99d735 100644 --- a/OpenCL/check_single_comp4_bs.c +++ b/OpenCL/check_single_comp4_bs.c @@ -1,3 +1,3 @@ const u32 final_hash_pos = digests_offset + 0; -mark_hash (plains_buf, d_result, salt_pos, 0, final_hash_pos, gid, il_pos + slice); +mark_hash (plains_buf, d_return_buf, salt_pos, 0, final_hash_pos, gid, il_pos + slice); diff --git a/OpenCL/m06211.cl b/OpenCL/m06211.cl index bece9b17d..a6d310250 100644 --- a/OpenCL/m06211.cl +++ b/OpenCL/m06211.cl @@ -670,7 +670,7 @@ __kernel void m06211_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -686,7 +686,7 @@ __kernel void m06211_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -702,7 +702,7 @@ __kernel void m06211_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } diff --git a/OpenCL/m06212.cl b/OpenCL/m06212.cl index 17293101c..0b7003f89 100644 --- a/OpenCL/m06212.cl +++ b/OpenCL/m06212.cl @@ -670,7 +670,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -686,7 +686,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -702,7 +702,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -741,7 +741,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -758,7 +758,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -775,7 +775,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } diff --git a/OpenCL/m06213.cl b/OpenCL/m06213.cl index 6a7910a4e..28cd4a624 100644 --- a/OpenCL/m06213.cl +++ b/OpenCL/m06213.cl @@ -670,7 +670,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -686,7 +686,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -702,7 +702,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -741,7 +741,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -758,7 +758,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -775,7 +775,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -815,7 +815,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -833,7 +833,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } diff --git a/OpenCL/m06221.cl b/OpenCL/m06221.cl index 7342f6579..477db802e 100644 --- a/OpenCL/m06221.cl +++ b/OpenCL/m06221.cl @@ -578,7 +578,7 @@ __kernel void m06221_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -594,7 +594,7 @@ __kernel void m06221_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -610,7 +610,7 @@ __kernel void m06221_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } diff --git a/OpenCL/m06222.cl b/OpenCL/m06222.cl index 6edee84aa..5b02b7478 100644 --- a/OpenCL/m06222.cl +++ b/OpenCL/m06222.cl @@ -578,7 +578,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -594,7 +594,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -610,7 +610,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -649,7 +649,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -666,7 +666,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -683,7 +683,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } diff --git a/OpenCL/m06223.cl b/OpenCL/m06223.cl index 1738da52f..92f30e0b0 100644 --- a/OpenCL/m06223.cl +++ b/OpenCL/m06223.cl @@ -578,7 +578,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -594,7 +594,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -610,7 +610,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -649,7 +649,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -666,7 +666,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -683,7 +683,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -723,7 +723,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -741,7 +741,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } diff --git a/OpenCL/m06231.cl b/OpenCL/m06231.cl index 39c001a8b..1f6fc824e 100644 --- a/OpenCL/m06231.cl +++ b/OpenCL/m06231.cl @@ -2215,7 +2215,7 @@ __kernel void m06231_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -2231,7 +2231,7 @@ __kernel void m06231_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -2247,7 +2247,7 @@ __kernel void m06231_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } diff --git a/OpenCL/m06232.cl b/OpenCL/m06232.cl index c5bc347ed..6a705ec3f 100644 --- a/OpenCL/m06232.cl +++ b/OpenCL/m06232.cl @@ -1984,7 +1984,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -2000,7 +2000,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -2016,7 +2016,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -2055,7 +2055,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -2072,7 +2072,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -2089,7 +2089,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } diff --git a/OpenCL/m06233.cl b/OpenCL/m06233.cl index 81180b9b0..2b858e9bb 100644 --- a/OpenCL/m06233.cl +++ b/OpenCL/m06233.cl @@ -1984,7 +1984,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -2000,7 +2000,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -2016,7 +2016,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -2055,7 +2055,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -2072,7 +2072,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -2089,7 +2089,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -2129,7 +2129,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } @@ -2147,7 +2147,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); d_return_buf[lid] = 1; } diff --git a/OpenCL/m06800.cl b/OpenCL/m06800.cl index 32d9a30be..cdcfed7b0 100644 --- a/OpenCL/m06800.cl +++ b/OpenCL/m06800.cl @@ -1573,7 +1573,7 @@ __kernel void m06800_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf && (out[2] == salt_buf[2]) && (out[3] == salt_buf[3])) { - mark_hash (plains_buf, hashes_shown, digests_offset + 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, 0); d_return_buf[lid] = 1; } diff --git a/OpenCL/m07500_a0.cl b/OpenCL/m07500_a0.cl index b3371297f..7f353f146 100644 --- a/OpenCL/m07500_a0.cl +++ b/OpenCL/m07500_a0.cl @@ -640,7 +640,7 @@ __kernel void m07500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1) { - mark_hash (plains_buf, hashes_shown, digests_offset, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos); d_return_buf[lid] = 1; } @@ -743,7 +743,7 @@ __kernel void m07500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1) { - mark_hash (plains_buf, hashes_shown, digests_offset, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos); d_return_buf[lid] = 1; } diff --git a/OpenCL/m07500_a1.cl b/OpenCL/m07500_a1.cl index c99bb8de7..a7ec17468 100644 --- a/OpenCL/m07500_a1.cl +++ b/OpenCL/m07500_a1.cl @@ -688,7 +688,7 @@ __kernel void m07500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1) { - mark_hash (plains_buf, hashes_shown, digests_offset, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos); d_return_buf[lid] = 1; } @@ -841,7 +841,7 @@ __kernel void m07500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1) { - mark_hash (plains_buf, hashes_shown, digests_offset, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos); d_return_buf[lid] = 1; } diff --git a/OpenCL/m07500_a3.cl b/OpenCL/m07500_a3.cl index 279ad1b8b..10d105ce4 100644 --- a/OpenCL/m07500_a3.cl +++ b/OpenCL/m07500_a3.cl @@ -634,7 +634,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, hashes_shown, digests_offset, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos); d_return_buf[lid] = 1; } diff --git a/OpenCL/m08800.cl b/OpenCL/m08800.cl index 934310bb8..c81dda6d6 100644 --- a/OpenCL/m08800.cl +++ b/OpenCL/m08800.cl @@ -1893,7 +1893,7 @@ __kernel void m08800_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf // MSDOS5.0 if ((r0 == 0x4f44534d) && (r1 == 0x302e3553)) { - mark_hash (plains_buf, hashes_shown, digests_offset + 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, 0); d_return_buf[lid] = 1; } @@ -1961,7 +1961,7 @@ __kernel void m08800_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if ((r[5] < 2) && (r[6] < 16) && ((r[14] & 0xffff) == 0xEF53)) { - mark_hash (plains_buf, hashes_shown, digests_offset + 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, 0); d_return_buf[lid] = 1; } diff --git a/OpenCL/m11300.cl b/OpenCL/m11300.cl index 46af4c222..04b11c828 100644 --- a/OpenCL/m11300.cl +++ b/OpenCL/m11300.cl @@ -1349,7 +1349,7 @@ __kernel void m11300_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf && (out[2] == 0x10101010) && (out[3] == 0x10101010)) { - mark_hash (plains_buf, hashes_shown, digests_offset + 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, 0); d_return_buf[lid] = 1; } diff --git a/OpenCL/m11600.cl b/OpenCL/m11600.cl index c83bf047e..0c0deff95 100644 --- a/OpenCL/m11600.cl +++ b/OpenCL/m11600.cl @@ -2015,7 +2015,7 @@ __kernel void m11600_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if ((out[0] == 0) && (out[1] == 0) && (out[2] == 0) && (out[3] == 0)) { - mark_hash (plains_buf, hashes_shown, digests_offset + 0, gid, 0); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, 0); d_return_buf[lid] = 1; } diff --git a/OpenCL/m13100_a0.cl b/OpenCL/m13100_a0.cl index 430e3d950..d36b4a399 100644 --- a/OpenCL/m13100_a0.cl +++ b/OpenCL/m13100_a0.cl @@ -779,7 +779,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, hashes_shown, digests_offset, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos); d_return_buf[lid] = 1; } @@ -866,7 +866,7 @@ __kernel void m13100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, 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, hashes_shown, digests_offset, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos); d_return_buf[lid] = 1; } @@ -961,7 +961,7 @@ __kernel void m13100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, 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, hashes_shown, digests_offset, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos); d_return_buf[lid] = 1; } diff --git a/OpenCL/m13100_a1.cl b/OpenCL/m13100_a1.cl index fc7840ab4..fdaf4adc4 100644 --- a/OpenCL/m13100_a1.cl +++ b/OpenCL/m13100_a1.cl @@ -859,7 +859,7 @@ __kernel void m13100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, 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, hashes_shown, digests_offset, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos); d_return_buf[lid] = 1; } @@ -1003,7 +1003,7 @@ __kernel void m13100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, 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, hashes_shown, digests_offset, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos); d_return_buf[lid] = 1; } diff --git a/OpenCL/m13100_a3.cl b/OpenCL/m13100_a3.cl index 4a2878e36..63a1e19d8 100644 --- a/OpenCL/m13100_a3.cl +++ b/OpenCL/m13100_a3.cl @@ -781,7 +781,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, hashes_shown, digests_offset, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos); d_return_buf[lid] = 1; } diff --git a/OpenCL/m13200.cl b/OpenCL/m13200.cl index 3daa846c7..d5e060653 100644 --- a/OpenCL/m13200.cl +++ b/OpenCL/m13200.cl @@ -1227,7 +1227,7 @@ __kernel void m13200_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if(tmps[gid].cipher[0]==0xA6A6A6A6 && tmps[gid].cipher[1]==0xA6A6A6A6) { - mark_hash (plains_buf, hashes_shown, digests_offset, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos); d_return_buf[lid] = 1; } diff --git a/OpenCL/m13400.cl b/OpenCL/m13400.cl index 1ba5d64c4..baf447030 100644 --- a/OpenCL/m13400.cl +++ b/OpenCL/m13400.cl @@ -1758,7 +1758,7 @@ __kernel void m13400_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf && esalt_bufs[salt_pos].contents_hash[6] == final_digest[6] && esalt_bufs[salt_pos].contents_hash[7] == final_digest[7]) { - mark_hash (plains_buf, hashes_shown, digests_offset, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos); d_return_buf[lid] = 1; } @@ -1964,7 +1964,7 @@ __kernel void m13400_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf && esalt_bufs[salt_pos].contents_hash[6] == final_digest[6] && esalt_bufs[salt_pos].contents_hash[7] == final_digest[7]) { - mark_hash (plains_buf, hashes_shown, digests_offset, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos); d_return_buf[lid] = 1; } @@ -2004,7 +2004,7 @@ __kernel void m13400_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf && esalt_bufs[salt_pos].expected_bytes[2] == out[2] && esalt_bufs[salt_pos].expected_bytes[3] == out[3]) { - mark_hash (plains_buf, hashes_shown, digests_offset, gid, il_pos); + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos); d_return_buf[lid] = 1; } diff --git a/OpenCL/simd.c b/OpenCL/simd.c index 829350db3..22a9f29ac 100644 --- a/OpenCL/simd.c +++ b/OpenCL/simd.c @@ -20,7 +20,7 @@ \ if (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos); \ } \ } \ } @@ -36,15 +36,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp0, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp0, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos); \ } \ } \ } \ @@ -69,7 +69,7 @@ \ if (vector_accessible (il_pos, il_cnt, 0) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 0); \ } \ } \ \ @@ -79,7 +79,7 @@ \ if (vector_accessible (il_pos, il_cnt, 1) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 1); \ } \ } \ } @@ -96,15 +96,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp0, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp0, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 0) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 0); \ } \ } \ } \ @@ -116,15 +116,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp1, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp1, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 1) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 1); \ } \ } \ } \ @@ -147,7 +147,7 @@ \ if (vector_accessible (il_pos, il_cnt, 0) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 0); \ } \ } \ \ @@ -157,7 +157,7 @@ \ if (vector_accessible (il_pos, il_cnt, 1) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 1); \ } \ } \ \ @@ -167,7 +167,7 @@ \ if (vector_accessible (il_pos, il_cnt, 2) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 2); \ } \ } \ \ @@ -177,7 +177,7 @@ \ if (vector_accessible (il_pos, il_cnt, 3) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 3); \ } \ } \ } @@ -196,15 +196,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp0, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp0, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 0) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 0); \ } \ } \ } \ @@ -216,15 +216,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp1, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp1, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 1) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 1); \ } \ } \ } \ @@ -236,15 +236,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp2, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp2, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 2) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 2); \ } \ } \ } \ @@ -256,15 +256,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp3, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp3, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 3) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 3); \ } \ } \ } \ @@ -287,7 +287,7 @@ \ if (vector_accessible (il_pos, il_cnt, 0) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 0); \ } \ } \ \ @@ -297,7 +297,7 @@ \ if (vector_accessible (il_pos, il_cnt, 1) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 1); \ } \ } \ \ @@ -307,7 +307,7 @@ \ if (vector_accessible (il_pos, il_cnt, 2) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 2); \ } \ } \ \ @@ -317,7 +317,7 @@ \ if (vector_accessible (il_pos, il_cnt, 3) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 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])) \ @@ -326,7 +326,7 @@ \ if (vector_accessible (il_pos, il_cnt, 4) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 4); \ } \ } \ \ @@ -336,7 +336,7 @@ \ if (vector_accessible (il_pos, il_cnt, 5) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 5); \ } \ } \ \ @@ -346,7 +346,7 @@ \ if (vector_accessible (il_pos, il_cnt, 6) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 6); \ } \ } \ \ @@ -356,7 +356,7 @@ \ if (vector_accessible (il_pos, il_cnt, 7) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 7); \ } \ } \ } @@ -379,15 +379,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp0, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp0, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 0) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 0); \ } \ } \ } \ @@ -399,15 +399,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp1, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp1, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 1) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 1); \ } \ } \ } \ @@ -419,15 +419,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp2, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp2, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 2) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 2); \ } \ } \ } \ @@ -439,15 +439,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp3, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp3, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 3) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 3); \ } \ } \ } \ @@ -458,15 +458,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp4, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp4, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 4) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 4); \ } \ } \ } \ @@ -478,15 +478,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp5, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp5, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 5) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 5); \ } \ } \ } \ @@ -498,15 +498,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp6, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp6, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 6) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 6); \ } \ } \ } \ @@ -518,15 +518,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp7, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp7, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 7) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 7); \ } \ } \ } \ @@ -549,7 +549,7 @@ \ if (vector_accessible (il_pos, il_cnt, 0) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 0); \ } \ } \ \ @@ -559,7 +559,7 @@ \ if (vector_accessible (il_pos, il_cnt, 1) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 1); \ } \ } \ \ @@ -569,7 +569,7 @@ \ if (vector_accessible (il_pos, il_cnt, 2) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 2); \ } \ } \ \ @@ -579,7 +579,7 @@ \ if (vector_accessible (il_pos, il_cnt, 3) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 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])) \ @@ -588,7 +588,7 @@ \ if (vector_accessible (il_pos, il_cnt, 4) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 4); \ } \ } \ \ @@ -598,7 +598,7 @@ \ if (vector_accessible (il_pos, il_cnt, 5) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 5); \ } \ } \ \ @@ -608,7 +608,7 @@ \ if (vector_accessible (il_pos, il_cnt, 6) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 6); \ } \ } \ \ @@ -618,7 +618,7 @@ \ if (vector_accessible (il_pos, il_cnt, 7) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 7); \ } \ } \ \ @@ -628,7 +628,7 @@ \ if (vector_accessible (il_pos, il_cnt, 8) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 8); \ } \ } \ \ @@ -638,7 +638,7 @@ \ if (vector_accessible (il_pos, il_cnt, 9) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 9); \ } \ } \ \ @@ -648,7 +648,7 @@ \ if (vector_accessible (il_pos, il_cnt, 10) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 10); \ } \ } \ \ @@ -658,7 +658,7 @@ \ if (vector_accessible (il_pos, il_cnt, 11) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 11); \ } \ } \ \ @@ -668,7 +668,7 @@ \ if (vector_accessible (il_pos, il_cnt, 12) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 12); \ } \ } \ \ @@ -678,7 +678,7 @@ \ if (vector_accessible (il_pos, il_cnt, 13) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 13); \ } \ } \ \ @@ -688,7 +688,7 @@ \ if (vector_accessible (il_pos, il_cnt, 14) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 14); \ } \ } \ \ @@ -698,7 +698,7 @@ \ if (vector_accessible (il_pos, il_cnt, 15) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, 0, final_hash_pos, gid, il_pos + 15); \ } \ } \ } @@ -729,15 +729,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp00, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp00, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 0) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 0); \ } \ } \ } \ @@ -749,15 +749,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp01, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp01, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 1) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 1); \ } \ } \ } \ @@ -769,15 +769,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp02, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp02, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 2) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 2); \ } \ } \ } \ @@ -789,15 +789,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp03, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp03, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 3) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 3); \ } \ } \ } \ @@ -809,15 +809,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp04, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp04, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 4) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 4); \ } \ } \ } \ @@ -829,15 +829,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp05, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp05, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 5) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 5); \ } \ } \ } \ @@ -849,15 +849,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp06, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp06, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 6) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 6); \ } \ } \ } \ @@ -869,15 +869,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp07, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp07, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 7) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 7); \ } \ } \ } \ @@ -889,15 +889,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp08, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp08, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 8) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 8); \ } \ } \ } \ @@ -909,15 +909,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp09, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp09, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 9) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 9); \ } \ } \ } \ @@ -929,15 +929,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp10, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp10, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 10) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 10); \ } \ } \ } \ @@ -949,15 +949,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp11, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp11, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 11) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 11); \ } \ } \ } \ @@ -969,15 +969,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp12, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp12, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 12) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 12); \ } \ } \ } \ @@ -989,15 +989,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp13, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp13, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 13) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 13); \ } \ } \ } \ @@ -1009,15 +1009,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp14, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp14, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 14) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 14); \ } \ } \ } \ @@ -1029,15 +1029,15 @@ bitmap_shift1, \ bitmap_shift2)) \ { \ - int digest_pos = find_hash (digest_tp15, digests_cnt, &digests_buf[digests_offset]); \ + int digest_pos = find_hash (digest_tp15, digests_cnt, &digests_buf[digests_offset]); \ \ - if (digest_pos != -1) \ + if (digest_pos != -1) \ { \ - const u32 final_hash_pos = digests_offset + digest_pos; \ + const u32 final_hash_pos = digests_offset + digest_pos; \ \ if (vector_accessible (il_pos, il_cnt, 15) && (atomic_add (&hashes_shown[final_hash_pos], 1) == 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, digest_pos, final_hash_pos, gid, il_pos + 15); \ } \ } \ } \