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

@@ -689,17 +689,26 @@ __kernel void m06211_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey2[7] = tmps[gid].out[15];
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    if (atomic_inc (&hashes_shown[0]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m06212.cl

@@ -697,19 +697,28 @@ __kernel void m06212_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey2[7] = tmps[gid].out[15];
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    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)
   volatile u32 ukey3[8];
@@ -742,17 +751,26 @@ __kernel void m06212_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey4[7] = tmps[gid].out[31];
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    if (atomic_inc (&hashes_shown[0]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m06213.cl

@@ -697,19 +697,28 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey2[7] = tmps[gid].out[15];
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    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)
   volatile u32 ukey3[8];
@@ -742,19 +751,28 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey4[7] = tmps[gid].out[31];
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    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)
   volatile u32 ukey5[8];
@@ -787,12 +805,18 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey6[7] = tmps[gid].out[47];
 
   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)
+  {
+    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)
+  {
+    if (atomic_inc (&hashes_shown[0]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m06221.cl

@@ -597,17 +597,26 @@ __kernel void m06221_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey2[7] = swap32 (l32_from_64 (tmps[gid].out[ 7]));
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    if (atomic_inc (&hashes_shown[0]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m06222.cl

@@ -605,19 +605,28 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey2[7] = swap32 (l32_from_64 (tmps[gid].out[ 7]));
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    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)
   volatile u32 ukey3[8];
@@ -650,17 +659,26 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey4[7] = swap32 (l32_from_64 (tmps[gid].out[15]));
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    if (atomic_inc (&hashes_shown[0]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m06223.cl

@@ -654,19 +654,28 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey2[7] = swap32 (l32_from_64 (tmps[gid].out[ 7]));
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    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)
   volatile u32 ukey3[8];
@@ -699,19 +708,28 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey4[7] = swap32 (l32_from_64 (tmps[gid].out[15]));
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    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)
   volatile u32 ukey5[8];
@@ -744,12 +762,18 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey6[7] = swap32 (l32_from_64 (tmps[gid].out[23]));
 
   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)
+  {
+    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)
+  {
+    if (atomic_inc (&hashes_shown[0]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m06231.cl

@@ -2232,17 +2232,26 @@ __kernel void m06231_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey2[7] = swap32 (tmps[gid].out[15]);
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    if (atomic_inc (&hashes_shown[0]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m06232.cl

@@ -2009,19 +2009,28 @@ __kernel void m06232_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey2[7] = swap32 (tmps[gid].out[15]);
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    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)
   volatile u32 ukey3[8];
@@ -2054,17 +2063,26 @@ __kernel void m06232_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey4[7] = swap32 (tmps[gid].out[31]);
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    if (atomic_inc (&hashes_shown[0]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m06233.cl

@@ -2009,19 +2009,28 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey2[7] = swap32 (tmps[gid].out[15]);
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    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)
   volatile u32 ukey3[8];
@@ -2054,19 +2063,28 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey4[7] = swap32 (tmps[gid].out[31]);
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    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)
   volatile u32 ukey5[8];
@@ -2099,12 +2117,18 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey6[7] = swap32 (tmps[gid].out[47]);
 
   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)
+  {
+    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)
+  {
+    if (atomic_inc (&hashes_shown[0]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m06800.cl

@@ -1562,9 +1562,12 @@ __kernel void m06800_comp (__global pw_t *pws, __global const kernel_rule_t *rul
    && (out[1] == salt_buf[1])
    && (out[2] == salt_buf[2])
    && (out[3] == salt_buf[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, 0);
     }
+  }
 
   /**
    * offline mode

OpenCL/m07500_a0.cl

@@ -631,10 +631,13 @@ __kernel void m07500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
     tmp[3] = digest[3];
 
     if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1)
+    {
+      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);
       }
     }
+  }
 }
 
 __kernel void m07500_m08 (__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 krb5pa_t *krb5pa_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)
@@ -732,10 +735,13 @@ __kernel void m07500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
     tmp[3] = digest[3];
 
     if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1)
+    {
+      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);
       }
     }
+  }
 }
 
 __kernel void m07500_s08 (__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 krb5pa_t *krb5pa_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)

OpenCL/m07500_a1.cl

@@ -679,10 +679,13 @@ __kernel void m07500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
     tmp[3] = digest[3];
 
     if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1)
+    {
+      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);
       }
     }
+  }
 }
 
 __kernel void m07500_m08 (__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 krb5pa_t *krb5pa_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)
@@ -830,10 +833,13 @@ __kernel void m07500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
     tmp[3] = digest[3];
 
     if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1)
+    {
+      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);
       }
     }
+  }
 }
 
 __kernel void m07500_s08 (__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 krb5pa_t *krb5pa_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)

OpenCL/m07500_a3.cl

@@ -623,10 +623,13 @@ void m07500 (__local RC4_KEY *rc4_keys, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[
     tmp[3] = digest[3];
 
     if (decrypt_and_check (&rc4_keys[lid], tmp, timestamp_ct) == 1)
+    {
+      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);
       }
     }
+  }
 }
 
 __kernel void m07500_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 krb5pa_t *krb5pa_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)

OpenCL/m08800.cl

@@ -1882,10 +1882,13 @@ __kernel void m08800_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
     // MSDOS5.0
     if ((r0 == 0x4f44534d) && (r1 == 0x302e3553))
+    {
+      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);
       }
     }
+  }
 
   // check for extfs
 
@@ -1948,8 +1951,11 @@ __kernel void m08800_comp (__global pw_t *pws, __global const kernel_rule_t *rul
     // has EXT2_SUPER_MAGIC
 
     if ((r[5] < 2) && (r[6] < 16) && ((r[14] & 0xffff) == 0xEF53))
+    {
+      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);
       }
     }
+  }
 }

OpenCL/m11300.cl

@@ -1338,7 +1338,10 @@ __kernel void m11300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
    && (out[1] == 0x10101010)
    && (out[2] == 0x10101010)
    && (out[3] == 0x10101010))
+  {
+    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);
     }
+  }
 }

OpenCL/m11600.cl

@@ -850,9 +850,12 @@ __kernel void m11600_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   if (gid >= gid_max) return;
 
   if (seven_zip_hook[gid].hook_success == 1)
+  {
+    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;
   }
- }
+}

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)
 {
   /**
@@ -852,10 +800,13 @@ __kernel void m13100_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
     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)
+    {
+      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);
       }
     }
+  }
 }
 
 __kernel void m13100_m08 (__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)
@@ -945,10 +896,13 @@ __kernel void m13100_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
     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)
+    {
+      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);
       }
     }
+  }
 }
 
 __kernel void m13100_s08 (__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)

OpenCL/m13100_a1.cl

@@ -847,10 +847,13 @@ __kernel void m13100_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
     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)
+    {
+      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);
       }
     }
+  }
 }
 
 __kernel void m13100_m08 (__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)
@@ -989,10 +992,13 @@ __kernel void m13100_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
     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)
+    {
+      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);
       }
     }
+  }
 }
 
 __kernel void m13100_s08 (__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)

OpenCL/m13100_a3.cl

@@ -769,10 +769,13 @@ void m13100 (__local RC4_KEY *rc4_keys, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[
     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)
+    {
+      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);
       }
     }
+  }
 }
 
 __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)

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)
+  {
+    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);
     }
+  }
 }

OpenCL/m13400.cl

@@ -991,10 +991,13 @@ __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
     {
       /* Construct final AES key */
@@ -1195,11 +1198,14 @@ __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
   {
     /* Construct final AES key */
@@ -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);
       }
     }
+  }
 }

OpenCL/m13751.cl

@@ -660,17 +660,26 @@ __kernel void m13751_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey2[7] = swap32 (tmps[gid].out[15]);
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    if (atomic_inc (&hashes_shown[digests_offset]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m13752.cl

@@ -660,19 +660,28 @@ __kernel void m13752_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey2[7] = swap32 (tmps[gid].out[15]);
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    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];
 
@@ -697,17 +706,26 @@ __kernel void m13752_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey4[7] = swap32 (tmps[gid].out[31]);
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    if (atomic_inc (&hashes_shown[digests_offset]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m13753.cl

@@ -660,19 +660,28 @@ __kernel void m13753_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey2[7] = swap32 (tmps[gid].out[15]);
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    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];
 
@@ -697,19 +706,28 @@ __kernel void m13753_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey4[7] = swap32 (tmps[gid].out[31]);
 
   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)
+  {
+    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)
+  {
+    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)
+  {
+    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];
 
@@ -734,12 +752,18 @@ __kernel void m13753_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey6[7] = swap32 (tmps[gid].out[47]);
 
   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)
+  {
+    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)
+  {
+    if (atomic_inc (&hashes_shown[digests_offset]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m14611.cl

@@ -698,7 +698,10 @@ __kernel void m14611_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   const float entropy = get_entropy (pt_buf, 128);
 
   if (entropy < MAX_ENTROPY)
+  {
+    if (atomic_inc (&hashes_shown[digests_offset]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m14612.cl

@@ -645,7 +645,10 @@ __kernel void m14612_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   const float entropy = get_entropy (pt_buf, 128);
 
   if (entropy < MAX_ENTROPY)
+  {
+    if (atomic_inc (&hashes_shown[digests_offset]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m14613.cl

@@ -645,7 +645,10 @@ __kernel void m14613_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   const float entropy = get_entropy (pt_buf, 128);
 
   if (entropy < MAX_ENTROPY)
+  {
+    if (atomic_inc (&hashes_shown[digests_offset]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m14621.cl

@@ -689,7 +689,10 @@ __kernel void m14621_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   const float entropy = get_entropy (pt_buf, 128);
 
   if (entropy < MAX_ENTROPY)
+  {
+    if (atomic_inc (&hashes_shown[digests_offset]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m14622.cl

@@ -636,7 +636,10 @@ __kernel void m14622_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   const float entropy = get_entropy (pt_buf, 128);
 
   if (entropy < MAX_ENTROPY)
+  {
+    if (atomic_inc (&hashes_shown[digests_offset]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m14623.cl

@@ -636,7 +636,10 @@ __kernel void m14623_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   const float entropy = get_entropy (pt_buf, 128);
 
   if (entropy < MAX_ENTROPY)
+  {
+    if (atomic_inc (&hashes_shown[digests_offset]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m14631.cl

@@ -665,7 +665,10 @@ __kernel void m14631_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   const float entropy = get_entropy (pt_buf, 128);
 
   if (entropy < MAX_ENTROPY)
+  {
+    if (atomic_inc (&hashes_shown[digests_offset]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m14632.cl

@@ -612,7 +612,10 @@ __kernel void m14632_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   const float entropy = get_entropy (pt_buf, 128);
 
   if (entropy < MAX_ENTROPY)
+  {
+    if (atomic_inc (&hashes_shown[digests_offset]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m14633.cl

@@ -612,7 +612,10 @@ __kernel void m14633_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   const float entropy = get_entropy (pt_buf, 128);
 
   if (entropy < MAX_ENTROPY)
+  {
+    if (atomic_inc (&hashes_shown[digests_offset]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m14641.cl

@@ -853,7 +853,10 @@ __kernel void m14641_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   const float entropy = get_entropy (pt_buf, 128);
 
   if (entropy < MAX_ENTROPY)
+  {
+    if (atomic_inc (&hashes_shown[digests_offset]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m14642.cl

@@ -800,7 +800,10 @@ __kernel void m14642_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   const float entropy = get_entropy (pt_buf, 128);
 
   if (entropy < MAX_ENTROPY)
+  {
+    if (atomic_inc (&hashes_shown[digests_offset]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m14643.cl

@@ -800,7 +800,10 @@ __kernel void m14643_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   const float entropy = get_entropy (pt_buf, 128);
 
   if (entropy < MAX_ENTROPY)
+  {
+    if (atomic_inc (&hashes_shown[digests_offset]) == 0)
     {
       mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
     }
+  }
 }

OpenCL/m14700.cl

@@ -1739,8 +1739,11 @@ __kernel void m14700_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 
   if ((cipher[0] == 0xa6a6a6a6) && (cipher[1] == 0xa6a6a6a6))
+  {
+    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;
   }

OpenCL/m14800.cl

@@ -2296,8 +2296,11 @@ __kernel void m14800_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 
   if ((cipher[0] == 0xa6a6a6a6) && (cipher[1] == 0xa6a6a6a6))
+  {
+    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;
   }

OpenCL/m15300.cl

@@ -2375,10 +2375,13 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
       &&  expectedHmac[2] == digest[2]
       &&  expectedHmac[3] == digest[3]
       &&  expectedHmac[4] == digest[4])
+    {
+      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)
   {
     /**
@@ -2627,8 +2630,11 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
        && expectedHmac[13] == l32_from_64_S (dgst64[6])
        && expectedHmac[14] == h32_from_64_S (dgst64[7])
        && expectedHmac[15] == l32_from_64_S (dgst64[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);
       }
     }
+  }
 }