Fix slow-hash-simd exits in _loop kernels

8 years ago · 8d1809cacb
parent 2d5c65fb4b
commit 8d1809cacb
7 changed files with 7 additions and 7 deletions
--- 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];
--- 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];
--- 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];
--- 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];
--- 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];
--- 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];
--- 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];