projects
/
hashcat.git
/ commitdiff
commit
grep
author
committer
pickaxe
?
search:
re
summary
|
shortlog
|
log
|
commit
| commitdiff |
tree
raw
|
patch
|
inline
| side by side (parent:
37953cd
)
Some fixes for rare kernels
author
jsteube
<jens.steube@gmail.com>
Fri, 20 May 2016 17:05:54 +0000
(19:05 +0200)
committer
jsteube
<jens.steube@gmail.com>
Fri, 20 May 2016 17:05:54 +0000
(19:05 +0200)
26 files changed:
OpenCL/check_multi_comp4.c
patch
|
blob
|
history
OpenCL/check_multi_comp4_bs.c
patch
|
blob
|
history
OpenCL/check_single_comp4.c
patch
|
blob
|
history
OpenCL/check_single_comp4_bs.c
patch
|
blob
|
history
OpenCL/m06211.cl
patch
|
blob
|
history
OpenCL/m06212.cl
patch
|
blob
|
history
OpenCL/m06213.cl
patch
|
blob
|
history
OpenCL/m06221.cl
patch
|
blob
|
history
OpenCL/m06222.cl
patch
|
blob
|
history
OpenCL/m06223.cl
patch
|
blob
|
history
OpenCL/m06231.cl
patch
|
blob
|
history
OpenCL/m06232.cl
patch
|
blob
|
history
OpenCL/m06233.cl
patch
|
blob
|
history
OpenCL/m06800.cl
patch
|
blob
|
history
OpenCL/m07500_a0.cl
patch
|
blob
|
history
OpenCL/m07500_a1.cl
patch
|
blob
|
history
OpenCL/m07500_a3.cl
patch
|
blob
|
history
OpenCL/m08800.cl
patch
|
blob
|
history
OpenCL/m11300.cl
patch
|
blob
|
history
OpenCL/m11600.cl
patch
|
blob
|
history
OpenCL/m13100_a0.cl
patch
|
blob
|
history
OpenCL/m13100_a1.cl
patch
|
blob
|
history
OpenCL/m13100_a3.cl
patch
|
blob
|
history
OpenCL/m13200.cl
patch
|
blob
|
history
OpenCL/m13400.cl
patch
|
blob
|
history
OpenCL/simd.c
patch
|
blob
|
history
diff --git
a/OpenCL/check_multi_comp4.c
b/OpenCL/check_multi_comp4.c
index
6b38416
..
f45a03b
100644
(file)
--- 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)
{
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
a2b371d
..
0fbcd7b
100644
(file)
--- 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)
{
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
7eba5e2
..
f754d8e
100644
(file)
--- 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)
{
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
f2c5366
..
6ec99d7
100644
(file)
--- 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;
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
bece9b1
..
a6d3102
100644
(file)
--- 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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m06212.cl
b/OpenCL/m06212.cl
index
1729310
..
0b7003f
100644
(file)
--- 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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m06213.cl
b/OpenCL/m06213.cl
index
6a7910a
..
28cd4a6
100644
(file)
--- 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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m06221.cl
b/OpenCL/m06221.cl
index
7342f65
..
477db80
100644
(file)
--- 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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m06222.cl
b/OpenCL/m06222.cl
index
6edee84
..
5b02b74
100644
(file)
--- 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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m06223.cl
b/OpenCL/m06223.cl
index
1738da5
..
92f30e0
100644
(file)
--- 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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m06231.cl
b/OpenCL/m06231.cl
index
39c001a
..
1f6fc82
100644
(file)
--- 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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m06232.cl
b/OpenCL/m06232.cl
index
c5bc347
..
6a705ec
100644
(file)
--- 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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m06233.cl
b/OpenCL/m06233.cl
index
81180b9
..
2b858e9
100644
(file)
--- 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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
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)))
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m06800.cl
b/OpenCL/m06800.cl
index
32d9a30
..
cdcfed7
100644
(file)
--- 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]))
{
&& (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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m07500_a0.cl
b/OpenCL/m07500_a0.cl
index
b337129
..
7f353f1
100644
(file)
--- 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)
{
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;
}
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)
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m07500_a1.cl
b/OpenCL/m07500_a1.cl
index
c99bb8d
..
a7ec174
100644
(file)
--- 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)
{
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;
}
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)
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m07500_a3.cl
b/OpenCL/m07500_a3.cl
index
279ad1b
..
10d105c
100644
(file)
--- 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)
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m08800.cl
b/OpenCL/m08800.cl
index
934310b
..
c81dda6
100644
(file)
--- 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))
{
// 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;
}
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))
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m11300.cl
b/OpenCL/m11300.cl
index
46af4c2
..
04b11c8
100644
(file)
--- 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))
{
&& (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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m11600.cl
b/OpenCL/m11600.cl
index
c83bf04
..
0c0deff
100644
(file)
--- 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))
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m13100_a0.cl
b/OpenCL/m13100_a0.cl
index
430e3d9
..
d36b4a3
100644
(file)
--- 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)
{
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;
}
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)
{
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;
}
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)
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m13100_a1.cl
b/OpenCL/m13100_a1.cl
index
fc7840a
..
fdaf4ad
100644
(file)
--- 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)
{
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;
}
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)
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m13100_a3.cl
b/OpenCL/m13100_a3.cl
index
4a2878e
..
63a1e19
100644
(file)
--- 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)
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m13200.cl
b/OpenCL/m13200.cl
index
3daa846
..
d5e0606
100644
(file)
--- 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)
{
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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/m13400.cl
b/OpenCL/m13400.cl
index
1ba5d64
..
baf4470
100644
(file)
--- 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])
{
&& 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;
}
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])
{
&& 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;
}
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])
{
&& 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;
}
d_return_buf[lid] = 1;
}
diff --git
a/OpenCL/simd.c
b/OpenCL/simd.c
index
829350d
..
22a9f29
100644
(file)
--- a/
OpenCL/simd.c
+++ b/
OpenCL/simd.c
@@
-20,7
+20,7
@@
\
if (atomic_add (&hashes_shown[final_hash_pos], 1) == 0) \
{ \
\
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)) \
{ \
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) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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])) \
} \
} \
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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])) \
} \
} \
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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)) \
{ \
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)) \
{ \
\
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); \
} \
} \
} \
} \
} \
} \