From cf57365e7c695f86f1e0fb1e0c7fe1790d038ba0 Mon Sep 17 00:00:00 2001 From: jsteube Date: Fri, 16 Jun 2017 10:48:10 +0200 Subject: [PATCH] Check hashes_shown[] whenever calling mark_hash directly. Not really needed right now (because those algorithms to have a single digest per salt), but this can change in the future --- OpenCL/m06211.cl | 15 ++++++++--- OpenCL/m06212.cl | 30 +++++++++++++++++----- OpenCL/m06213.cl | 40 +++++++++++++++++++++++------ OpenCL/m06221.cl | 15 ++++++++--- OpenCL/m06222.cl | 30 +++++++++++++++++----- OpenCL/m06223.cl | 40 +++++++++++++++++++++++------ OpenCL/m06231.cl | 15 ++++++++--- OpenCL/m06232.cl | 30 +++++++++++++++++----- OpenCL/m06233.cl | 40 +++++++++++++++++++++++------ OpenCL/m06800.cl | 5 +++- OpenCL/m07500_a0.cl | 10 ++++++-- OpenCL/m07500_a1.cl | 10 ++++++-- OpenCL/m07500_a3.cl | 5 +++- OpenCL/m08800.cl | 10 ++++++-- OpenCL/m11300.cl | 5 +++- OpenCL/m11600.cl | 7 +++-- OpenCL/m13100_a0.cl | 62 ++++++--------------------------------------- OpenCL/m13100_a1.cl | 10 ++++++-- OpenCL/m13100_a3.cl | 5 +++- OpenCL/m13200.cl | 7 +++-- OpenCL/m13400.cl | 9 +++++++ OpenCL/m13751.cl | 15 ++++++++--- OpenCL/m13752.cl | 30 +++++++++++++++++----- OpenCL/m13753.cl | 40 +++++++++++++++++++++++------ OpenCL/m14611.cl | 5 +++- OpenCL/m14612.cl | 5 +++- OpenCL/m14613.cl | 5 +++- OpenCL/m14621.cl | 5 +++- OpenCL/m14622.cl | 5 +++- OpenCL/m14623.cl | 5 +++- OpenCL/m14631.cl | 5 +++- OpenCL/m14632.cl | 5 +++- OpenCL/m14633.cl | 5 +++- OpenCL/m14641.cl | 5 +++- OpenCL/m14642.cl | 5 +++- OpenCL/m14643.cl | 5 +++- OpenCL/m14700.cl | 5 +++- OpenCL/m14800.cl | 5 +++- OpenCL/m15300.cl | 10 ++++++-- 39 files changed, 411 insertions(+), 154 deletions(-) diff --git a/OpenCL/m06211.cl b/OpenCL/m06211.cl index 034a7befb..65c165338 100644 --- a/OpenCL/m06211.cl +++ b/OpenCL/m06211.cl @@ -690,16 +690,25 @@ __kernel void m06211_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m06212.cl b/OpenCL/m06212.cl index 7c4022081..e25740296 100644 --- a/OpenCL/m06212.cl +++ b/OpenCL/m06212.cl @@ -698,17 +698,26 @@ __kernel void m06212_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } #if defined (IS_APPLE) && defined (IS_GPU) @@ -743,16 +752,25 @@ __kernel void m06212_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m06213.cl b/OpenCL/m06213.cl index c00adcd6d..b45b4adc4 100644 --- a/OpenCL/m06213.cl +++ b/OpenCL/m06213.cl @@ -698,17 +698,26 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } #if defined (IS_APPLE) && defined (IS_GPU) @@ -743,17 +752,26 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } #if defined (IS_APPLE) && defined (IS_GPU) @@ -788,11 +806,17 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent_twofish_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m06221.cl b/OpenCL/m06221.cl index 1609b0b86..ece2fcc74 100644 --- a/OpenCL/m06221.cl +++ b/OpenCL/m06221.cl @@ -598,16 +598,25 @@ __kernel void m06221_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m06222.cl b/OpenCL/m06222.cl index fddf24dce..96c2b37ac 100644 --- a/OpenCL/m06222.cl +++ b/OpenCL/m06222.cl @@ -606,17 +606,26 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } #if defined (IS_APPLE) && defined (IS_GPU) @@ -651,16 +660,25 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m06223.cl b/OpenCL/m06223.cl index d55d58075..bba68d13c 100644 --- a/OpenCL/m06223.cl +++ b/OpenCL/m06223.cl @@ -655,17 +655,26 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } #if defined (IS_APPLE) && defined (IS_GPU) @@ -700,17 +709,26 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } #if defined (IS_APPLE) && defined (IS_GPU) @@ -745,11 +763,17 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent_twofish_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m06231.cl b/OpenCL/m06231.cl index 2841c421d..418fc41d3 100644 --- a/OpenCL/m06231.cl +++ b/OpenCL/m06231.cl @@ -2233,16 +2233,25 @@ __kernel void m06231_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m06232.cl b/OpenCL/m06232.cl index 883b54963..bb72e50be 100644 --- a/OpenCL/m06232.cl +++ b/OpenCL/m06232.cl @@ -2010,17 +2010,26 @@ __kernel void m06232_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } #if defined (IS_APPLE) && defined (IS_GPU) @@ -2055,16 +2064,25 @@ __kernel void m06232_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m06233.cl b/OpenCL/m06233.cl index d4a539816..c3e06b865 100644 --- a/OpenCL/m06233.cl +++ b/OpenCL/m06233.cl @@ -2010,17 +2010,26 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } #if defined (IS_APPLE) && defined (IS_GPU) @@ -2055,17 +2064,26 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } #if defined (IS_APPLE) && defined (IS_GPU) @@ -2100,11 +2118,17 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent_twofish_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[0]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m06800.cl b/OpenCL/m06800.cl index e2c90eda5..96a95a5c9 100644 --- a/OpenCL/m06800.cl +++ b/OpenCL/m06800.cl @@ -1563,7 +1563,10 @@ __kernel void m06800_comp (__global pw_t *pws, __global const kernel_rule_t *rul && (out[2] == salt_buf[2]) && (out[3] == salt_buf[3])) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0); + } } /** diff --git a/OpenCL/m07500_a0.cl b/OpenCL/m07500_a0.cl index 2bc72f095..f624466b3 100644 --- a/OpenCL/m07500_a0.cl +++ b/OpenCL/m07500_a0.cl @@ -632,7 +632,10 @@ __kernel void m07500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + } } } } @@ -733,7 +736,10 @@ __kernel void m07500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + } } } } diff --git a/OpenCL/m07500_a1.cl b/OpenCL/m07500_a1.cl index d9cc40cc6..829ec49ba 100644 --- a/OpenCL/m07500_a1.cl +++ b/OpenCL/m07500_a1.cl @@ -680,7 +680,10 @@ __kernel void m07500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + } } } } @@ -831,7 +834,10 @@ __kernel void m07500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + } } } } diff --git a/OpenCL/m07500_a3.cl b/OpenCL/m07500_a3.cl index 26e7951d4..77084b9dd 100644 --- a/OpenCL/m07500_a3.cl +++ b/OpenCL/m07500_a3.cl @@ -624,7 +624,10 @@ void m07500 (__local RC4_KEY *rc4_keys, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[ if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + } } } } diff --git a/OpenCL/m08800.cl b/OpenCL/m08800.cl index ac3fb46ce..2c025ebbd 100644 --- a/OpenCL/m08800.cl +++ b/OpenCL/m08800.cl @@ -1883,7 +1883,10 @@ __kernel void m08800_comp (__global pw_t *pws, __global const kernel_rule_t *rul // MSDOS5.0 if ((r0 == 0x4f44534d) && (r1 == 0x302e3553)) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0); + } } } @@ -1949,7 +1952,10 @@ __kernel void m08800_comp (__global pw_t *pws, __global const kernel_rule_t *rul if ((r[5] < 2) && (r[6] < 16) && ((r[14] & 0xffff) == 0xEF53)) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0); + } } } } diff --git a/OpenCL/m11300.cl b/OpenCL/m11300.cl index dc07eace3..38fdabb28 100644 --- a/OpenCL/m11300.cl +++ b/OpenCL/m11300.cl @@ -1339,6 +1339,9 @@ __kernel void m11300_comp (__global pw_t *pws, __global const kernel_rule_t *rul && (out[2] == 0x10101010) && (out[3] == 0x10101010)) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0); + } } } diff --git a/OpenCL/m11600.cl b/OpenCL/m11600.cl index 440afede1..9d6eebe64 100644 --- a/OpenCL/m11600.cl +++ b/OpenCL/m11600.cl @@ -851,8 +851,11 @@ __kernel void m11600_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (seven_zip_hook[gid].hook_success == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0); + } return; } - } +} diff --git a/OpenCL/m13100_a0.cl b/OpenCL/m13100_a0.cl index 08f4dd70c..f9c214baa 100644 --- a/OpenCL/m13100_a0.cl +++ b/OpenCL/m13100_a0.cl @@ -721,58 +721,6 @@ void kerb_prepare (const u32 w0[4], const u32 w1[4], const u32 pw_len, const u32 hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest); } -void m13100 (__local RC4_KEY *rc4_keys, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) -{ - /** - * modifier - */ - - const u32 gid = get_global_id (0); - const u32 lid = get_local_id (0); - - /** - * salt - */ - - u32 checksum[4]; - - checksum[0] = krb5tgs_bufs[digests_offset].checksum[0]; - checksum[1] = krb5tgs_bufs[digests_offset].checksum[1]; - checksum[2] = krb5tgs_bufs[digests_offset].checksum[2]; - checksum[3] = krb5tgs_bufs[digests_offset].checksum[3]; - - /** - * loop - */ - - u32 w0l = w0[0]; - - for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) - { - const u32 w0r = bfs_buf[il_pos].i; - - w0[0] = w0l | w0r; - - u32 digest[4]; - - u32 K2[4]; - - kerb_prepare (w0, w1, pw_len, checksum, digest, K2); - - u32 tmp[4]; - - tmp[0] = digest[0]; - tmp[1] = digest[1]; - tmp[2] = digest[2]; - tmp[3] = digest[3]; - - if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1) - { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); - } - } -} - __kernel void m13100_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** @@ -853,7 +801,10 @@ __kernel void m13100_m04 (__global pw_t *pws, __global const kernel_rule_t *rule if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + } } } } @@ -946,7 +897,10 @@ __kernel void m13100_s04 (__global pw_t *pws, __global const kernel_rule_t *rule if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + } } } } diff --git a/OpenCL/m13100_a1.cl b/OpenCL/m13100_a1.cl index 356a7dbd8..e89866cfa 100644 --- a/OpenCL/m13100_a1.cl +++ b/OpenCL/m13100_a1.cl @@ -848,7 +848,10 @@ __kernel void m13100_m04 (__global pw_t *pws, __global const kernel_rule_t *rule if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + } } } } @@ -990,7 +993,10 @@ __kernel void m13100_s04 (__global pw_t *pws, __global const kernel_rule_t *rule if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + } } } } diff --git a/OpenCL/m13100_a3.cl b/OpenCL/m13100_a3.cl index 62a4b9583..e5d33a4e4 100644 --- a/OpenCL/m13100_a3.cl +++ b/OpenCL/m13100_a3.cl @@ -770,7 +770,10 @@ 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[digests_offset].edata2, krb5tgs_bufs[digests_offset].edata2_len, K2, checksum) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + } } } } diff --git a/OpenCL/m13200.cl b/OpenCL/m13200.cl index d5c20fe57..3b0d96d91 100644 --- a/OpenCL/m13200.cl +++ b/OpenCL/m13200.cl @@ -1206,8 +1206,11 @@ __kernel void m13200_comp (__global pw_t *pws, __global const kernel_rule_t *rul #define il_pos 0 - 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, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + } } } diff --git a/OpenCL/m13400.cl b/OpenCL/m13400.cl index 78bb3dc09..ff31f459b 100644 --- a/OpenCL/m13400.cl +++ b/OpenCL/m13400.cl @@ -991,9 +991,12 @@ __kernel void m13400_comp (__global pw_t *pws, __global const kernel_rule_t *rul && esalt_bufs[digests_offset].contents_hash[5] == final_digest[5] && esalt_bufs[digests_offset].contents_hash[6] == final_digest[6] && esalt_bufs[digests_offset].contents_hash[7] == final_digest[7]) + { + if (atomic_inc (&hashes_shown[digests_offset]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); } + } } else { @@ -1195,9 +1198,12 @@ __kernel void m13400_comp (__global pw_t *pws, __global const kernel_rule_t *rul && esalt_bufs[digests_offset].contents_hash[5] == final_digest[5] && esalt_bufs[digests_offset].contents_hash[6] == final_digest[6] && esalt_bufs[digests_offset].contents_hash[7] == final_digest[7]) + { + if (atomic_inc (&hashes_shown[digests_offset]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); } + } } } else @@ -1233,8 +1239,11 @@ __kernel void m13400_comp (__global pw_t *pws, __global const kernel_rule_t *rul && esalt_bufs[digests_offset].expected_bytes[1] == out[1] && esalt_bufs[digests_offset].expected_bytes[2] == out[2] && esalt_bufs[digests_offset].expected_bytes[3] == out[3]) + { + if (atomic_inc (&hashes_shown[digests_offset]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); } + } } } diff --git a/OpenCL/m13751.cl b/OpenCL/m13751.cl index 3430253ff..4830c6dda 100644 --- a/OpenCL/m13751.cl +++ b/OpenCL/m13751.cl @@ -661,16 +661,25 @@ __kernel void m13751_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m13752.cl b/OpenCL/m13752.cl index 6c461f23b..ac8670556 100644 --- a/OpenCL/m13752.cl +++ b/OpenCL/m13752.cl @@ -661,17 +661,26 @@ __kernel void m13752_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } u32 ukey3[8]; @@ -698,16 +707,25 @@ __kernel void m13752_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m13753.cl b/OpenCL/m13753.cl index f81dcb90c..c2ca6debb 100644 --- a/OpenCL/m13753.cl +++ b/OpenCL/m13753.cl @@ -661,17 +661,26 @@ __kernel void m13753_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } u32 ukey3[8]; @@ -698,17 +707,26 @@ __kernel void m13753_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } u32 ukey5[8]; @@ -735,11 +753,17 @@ __kernel void m13753_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (verify_header_aes_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } if (verify_header_serpent_twofish_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m14611.cl b/OpenCL/m14611.cl index 4a093febe..6ff523a59 100644 --- a/OpenCL/m14611.cl +++ b/OpenCL/m14611.cl @@ -699,6 +699,9 @@ __kernel void m14611_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (entropy < MAX_ENTROPY) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m14612.cl b/OpenCL/m14612.cl index b24600c1e..dea8a8f4a 100644 --- a/OpenCL/m14612.cl +++ b/OpenCL/m14612.cl @@ -646,6 +646,9 @@ __kernel void m14612_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (entropy < MAX_ENTROPY) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m14613.cl b/OpenCL/m14613.cl index 1c0237364..b58a3b252 100644 --- a/OpenCL/m14613.cl +++ b/OpenCL/m14613.cl @@ -646,6 +646,9 @@ __kernel void m14613_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (entropy < MAX_ENTROPY) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m14621.cl b/OpenCL/m14621.cl index 128bfd5a1..05ba2f91f 100644 --- a/OpenCL/m14621.cl +++ b/OpenCL/m14621.cl @@ -690,6 +690,9 @@ __kernel void m14621_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (entropy < MAX_ENTROPY) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m14622.cl b/OpenCL/m14622.cl index 32971b0c6..83ef415ba 100644 --- a/OpenCL/m14622.cl +++ b/OpenCL/m14622.cl @@ -637,6 +637,9 @@ __kernel void m14622_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (entropy < MAX_ENTROPY) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m14623.cl b/OpenCL/m14623.cl index 114841683..c32b6bbb8 100644 --- a/OpenCL/m14623.cl +++ b/OpenCL/m14623.cl @@ -637,6 +637,9 @@ __kernel void m14623_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (entropy < MAX_ENTROPY) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m14631.cl b/OpenCL/m14631.cl index 7f59bb52c..742f14fd5 100644 --- a/OpenCL/m14631.cl +++ b/OpenCL/m14631.cl @@ -666,6 +666,9 @@ __kernel void m14631_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (entropy < MAX_ENTROPY) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m14632.cl b/OpenCL/m14632.cl index 965430f26..0e7c80ce1 100644 --- a/OpenCL/m14632.cl +++ b/OpenCL/m14632.cl @@ -613,6 +613,9 @@ __kernel void m14632_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (entropy < MAX_ENTROPY) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m14633.cl b/OpenCL/m14633.cl index 2ad798344..6a3a5bf15 100644 --- a/OpenCL/m14633.cl +++ b/OpenCL/m14633.cl @@ -613,6 +613,9 @@ __kernel void m14633_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (entropy < MAX_ENTROPY) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m14641.cl b/OpenCL/m14641.cl index 19c15df3d..90df7cf08 100644 --- a/OpenCL/m14641.cl +++ b/OpenCL/m14641.cl @@ -854,6 +854,9 @@ __kernel void m14641_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (entropy < MAX_ENTROPY) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m14642.cl b/OpenCL/m14642.cl index 6c2bd437c..c2ec92ec8 100644 --- a/OpenCL/m14642.cl +++ b/OpenCL/m14642.cl @@ -801,6 +801,9 @@ __kernel void m14642_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (entropy < MAX_ENTROPY) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m14643.cl b/OpenCL/m14643.cl index 230a1f951..07f44d6dd 100644 --- a/OpenCL/m14643.cl +++ b/OpenCL/m14643.cl @@ -801,6 +801,9 @@ __kernel void m14643_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (entropy < MAX_ENTROPY) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); + } } } diff --git a/OpenCL/m14700.cl b/OpenCL/m14700.cl index 4ab78b9a5..40c8bba4a 100644 --- a/OpenCL/m14700.cl +++ b/OpenCL/m14700.cl @@ -1740,7 +1740,10 @@ __kernel void m14700_comp (__global pw_t *pws, __global const kernel_rule_t *rul if ((cipher[0] == 0xa6a6a6a6) && (cipher[1] == 0xa6a6a6a6)) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0); + } return; } diff --git a/OpenCL/m14800.cl b/OpenCL/m14800.cl index aa2e58b3f..2802175cc 100644 --- a/OpenCL/m14800.cl +++ b/OpenCL/m14800.cl @@ -2297,7 +2297,10 @@ __kernel void m14800_comp (__global pw_t *pws, __global const kernel_rule_t *rul if ((cipher[0] == 0xa6a6a6a6) && (cipher[1] == 0xa6a6a6a6)) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, 0); + } return; } diff --git a/OpenCL/m15300.cl b/OpenCL/m15300.cl index 11a22ace1..4a63b36a4 100644 --- a/OpenCL/m15300.cl +++ b/OpenCL/m15300.cl @@ -2376,7 +2376,10 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul && expectedHmac[3] == digest[3] && expectedHmac[4] == digest[4]) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + } } } else if (esalt_bufs[digests_offset].version == 2) @@ -2628,7 +2631,10 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul && expectedHmac[14] == h32_from_64_S (dgst64[7]) && expectedHmac[15] == l32_from_64_S (dgst64[7])) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + if (atomic_inc (&hashes_shown[digests_offset]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + } } } }