From 8d1809cacb58a3681298e852dc66ce7904d3d9e9 Mon Sep 17 00:00:00 2001 From: jsteube Date: Sat, 29 Oct 2016 13:49:08 +0200 Subject: [PATCH] Fix slow-hash-simd exits in _loop kernels --- OpenCL/m00400.cl | 2 +- OpenCL/m02100.cl | 2 +- OpenCL/m02500.cl | 2 +- OpenCL/m07100.cl | 2 +- OpenCL/m10900.cl | 2 +- OpenCL/m11900.cl | 2 +- OpenCL/m12000.cl | 2 +- 7 files changed, 7 insertions(+), 7 deletions(-) diff --git a/OpenCL/m00400.cl b/OpenCL/m00400.cl index 3736c0c50..4cd755ffd 100644 --- a/OpenCL/m00400.cl +++ b/OpenCL/m00400.cl @@ -318,7 +318,7 @@ __kernel void m00400_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf const u32 gid = get_global_id (0); - if (gid >= gid_max) return; + if ((gid * VECT_SIZE) >= gid_max) return; u32x w0[4]; u32x w1[4]; diff --git a/OpenCL/m02100.cl b/OpenCL/m02100.cl index d3dc3c240..d72e6b0a9 100644 --- a/OpenCL/m02100.cl +++ b/OpenCL/m02100.cl @@ -671,7 +671,7 @@ __kernel void m02100_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf const u32 gid = get_global_id (0); - if (gid >= gid_max) return; + if ((gid * VECT_SIZE) >= gid_max) return; u32x ipad[5]; u32x opad[5]; diff --git a/OpenCL/m02500.cl b/OpenCL/m02500.cl index 170b847cc..6247605b3 100644 --- a/OpenCL/m02500.cl +++ b/OpenCL/m02500.cl @@ -740,7 +740,7 @@ __kernel void m02500_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf { const u32 gid = get_global_id (0); - if (gid >= gid_max) return; + if ((gid * VECT_SIZE) >= gid_max) return; u32x ipad[5]; u32x opad[5]; diff --git a/OpenCL/m07100.cl b/OpenCL/m07100.cl index 4f8e825bf..ad4219f43 100644 --- a/OpenCL/m07100.cl +++ b/OpenCL/m07100.cl @@ -611,7 +611,7 @@ __kernel void m07100_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf { const u32 gid = get_global_id (0); - if (gid >= gid_max) return; + if ((gid * VECT_SIZE) >= gid_max) return; u64x ipad[8]; u64x opad[8]; diff --git a/OpenCL/m10900.cl b/OpenCL/m10900.cl index af74df3c6..0320eec13 100644 --- a/OpenCL/m10900.cl +++ b/OpenCL/m10900.cl @@ -534,7 +534,7 @@ __kernel void m10900_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf { const u32 gid = get_global_id (0); - if (gid >= gid_max) return; + if ((gid * VECT_SIZE) >= gid_max) return; u32x ipad[8]; u32x opad[8]; diff --git a/OpenCL/m11900.cl b/OpenCL/m11900.cl index 054a573d6..95194173b 100644 --- a/OpenCL/m11900.cl +++ b/OpenCL/m11900.cl @@ -486,7 +486,7 @@ __kernel void m11900_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf { const u32 gid = get_global_id (0); - if (gid >= gid_max) return; + if ((gid * VECT_SIZE) >= gid_max) return; u32x ipad[4]; u32x opad[4]; diff --git a/OpenCL/m12000.cl b/OpenCL/m12000.cl index 71200d31d..9854f87ba 100644 --- a/OpenCL/m12000.cl +++ b/OpenCL/m12000.cl @@ -558,7 +558,7 @@ __kernel void m12000_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf { const u32 gid = get_global_id (0); - if (gid >= gid_max) return; + if ((gid * VECT_SIZE) >= gid_max) return; u32x ipad[5]; u32x opad[5];