From 0a676b549f1bdb5ae3a2d2d1d8e2bd36a9817451 Mon Sep 17 00:00:00 2001 From: jsteube Date: Mon, 7 Aug 2017 17:25:15 +0200 Subject: [PATCH] Remove global barrier when not needed to workaround Intel OpenCL runtime bug --- OpenCL/m00000_a0.cl | 4 ---- OpenCL/m00000_a3.cl | 4 ---- OpenCL/m00010_a0.cl | 8 -------- OpenCL/m00010_a1.cl | 4 ---- OpenCL/m00010_a3.cl | 8 -------- OpenCL/m00020_a0.cl | 4 ---- OpenCL/m00020_a3.cl | 4 ---- OpenCL/m00030_a0.cl | 8 -------- OpenCL/m00030_a1.cl | 4 ---- OpenCL/m00030_a3.cl | 8 -------- OpenCL/m00040_a0.cl | 4 ---- OpenCL/m00040_a3.cl | 4 ---- OpenCL/m00050_a0.cl | 8 -------- OpenCL/m00050_a1.cl | 8 -------- OpenCL/m00050_a3.cl | 8 -------- OpenCL/m00060_a0.cl | 8 -------- OpenCL/m00060_a1.cl | 8 -------- OpenCL/m00060_a3.cl | 8 -------- OpenCL/m00100_a0.cl | 4 ---- OpenCL/m00100_a3.cl | 4 ---- OpenCL/m00110_a0.cl | 8 -------- OpenCL/m00110_a1.cl | 4 ---- OpenCL/m00110_a3.cl | 8 -------- OpenCL/m00120_a0.cl | 4 ---- OpenCL/m00120_a3.cl | 4 ---- OpenCL/m00130_a0.cl | 8 -------- OpenCL/m00130_a1.cl | 4 ---- OpenCL/m00130_a3.cl | 8 -------- OpenCL/m00140_a0.cl | 4 ---- OpenCL/m00140_a3.cl | 4 ---- OpenCL/m00150_a0.cl | 8 -------- OpenCL/m00150_a1.cl | 8 -------- OpenCL/m00150_a3.cl | 8 -------- OpenCL/m00160_a0.cl | 8 -------- OpenCL/m00160_a1.cl | 8 -------- OpenCL/m00160_a3.cl | 8 -------- OpenCL/m00300_a0.cl | 4 ---- OpenCL/m00300_a3.cl | 4 ---- OpenCL/m00400.cl | 2 -- OpenCL/m00500.cl | 8 -------- OpenCL/m00900_a0.cl | 4 ---- OpenCL/m00900_a3.cl | 4 ---- OpenCL/m01000_a0.cl | 4 ---- OpenCL/m01000_a3.cl | 4 ---- OpenCL/m01100_a0.cl | 8 -------- OpenCL/m01100_a1.cl | 4 ---- OpenCL/m01100_a3.cl | 8 -------- OpenCL/m01300_a0.cl | 4 ---- OpenCL/m01300_a3.cl | 4 ---- OpenCL/m01400_a0.cl | 4 ---- OpenCL/m01400_a3.cl | 4 ---- OpenCL/m01410_a0.cl | 8 -------- OpenCL/m01410_a1.cl | 4 ---- OpenCL/m01410_a3.cl | 8 -------- OpenCL/m01420_a0.cl | 4 ---- OpenCL/m01420_a3.cl | 4 ---- OpenCL/m01430_a0.cl | 8 -------- OpenCL/m01430_a1.cl | 4 ---- OpenCL/m01430_a3.cl | 8 -------- OpenCL/m01440_a0.cl | 4 ---- OpenCL/m01440_a3.cl | 4 ---- OpenCL/m01450_a0.cl | 8 -------- OpenCL/m01450_a1.cl | 8 -------- OpenCL/m01450_a3.cl | 8 -------- OpenCL/m01460_a0.cl | 8 -------- OpenCL/m01460_a1.cl | 8 -------- OpenCL/m01460_a3.cl | 8 -------- OpenCL/m01600.cl | 8 -------- OpenCL/m01700_a0.cl | 4 ---- OpenCL/m01700_a3.cl | 4 ---- OpenCL/m01710_a0.cl | 8 -------- OpenCL/m01710_a1.cl | 4 ---- OpenCL/m01710_a3.cl | 8 -------- OpenCL/m01720_a0.cl | 4 ---- OpenCL/m01720_a3.cl | 4 ---- OpenCL/m01730_a0.cl | 8 -------- OpenCL/m01730_a1.cl | 4 ---- OpenCL/m01730_a3.cl | 8 -------- OpenCL/m01740_a0.cl | 4 ---- OpenCL/m01740_a3.cl | 4 ---- OpenCL/m01750_a0.cl | 8 -------- OpenCL/m01750_a1.cl | 8 -------- OpenCL/m01750_a3.cl | 8 -------- OpenCL/m01760_a0.cl | 8 -------- OpenCL/m01760_a1.cl | 8 -------- OpenCL/m01760_a3.cl | 8 -------- OpenCL/m01800.cl | 4 ---- OpenCL/m02610_a0.cl | 8 -------- OpenCL/m02610_a1.cl | 4 ---- OpenCL/m02610_a3.cl | 8 -------- OpenCL/m02810_a0.cl | 8 -------- OpenCL/m02810_a1.cl | 4 ---- OpenCL/m02810_a3.cl | 8 -------- OpenCL/m03710_a0.cl | 8 -------- OpenCL/m03710_a1.cl | 4 ---- OpenCL/m03710_a3.cl | 8 -------- OpenCL/m03800_a0.cl | 8 -------- OpenCL/m03800_a1.cl | 4 ---- OpenCL/m03800_a3.cl | 8 -------- OpenCL/m03910_a0.cl | 8 -------- OpenCL/m03910_a1.cl | 4 ---- OpenCL/m03910_a3.cl | 8 -------- OpenCL/m04010_a0.cl | 4 ---- OpenCL/m04010_a3.cl | 4 ---- OpenCL/m04110_a0.cl | 8 -------- OpenCL/m04110_a1.cl | 4 ---- OpenCL/m04110_a3.cl | 8 -------- OpenCL/m04310_a0.cl | 8 -------- OpenCL/m04310_a1.cl | 4 ---- OpenCL/m04310_a3.cl | 8 -------- OpenCL/m04400_a0.cl | 4 ---- OpenCL/m04400_a3.cl | 4 ---- OpenCL/m04500_a0.cl | 4 ---- OpenCL/m04500_a3.cl | 4 ---- OpenCL/m04520_a0.cl | 4 ---- OpenCL/m04520_a3.cl | 4 ---- OpenCL/m04700_a0.cl | 4 ---- OpenCL/m04700_a3.cl | 4 ---- OpenCL/m04800_a0.cl | 4 ---- OpenCL/m04800_a3.cl | 4 ---- OpenCL/m04900_a0.cl | 8 -------- OpenCL/m04900_a1.cl | 4 ---- OpenCL/m04900_a3.cl | 8 -------- OpenCL/m05100_a0.cl | 4 ---- OpenCL/m05100_a3.cl | 4 ---- OpenCL/m05300_a0.cl | 4 ---- OpenCL/m05300_a1.cl | 4 ---- OpenCL/m05300_a3.cl | 4 ---- OpenCL/m05400_a0.cl | 4 ---- OpenCL/m05400_a1.cl | 4 ---- OpenCL/m05400_a3.cl | 4 ---- OpenCL/m05500_a0.cl | 4 ---- OpenCL/m05500_a3.cl | 4 ---- OpenCL/m05600_a0.cl | 4 ---- OpenCL/m05600_a3.cl | 4 ---- OpenCL/m05800.cl | 4 ---- OpenCL/m06000_a0.cl | 4 ---- OpenCL/m06000_a3.cl | 4 ---- OpenCL/m06100_a0.cl | 4 ---- OpenCL/m06100_a3.cl | 4 ---- OpenCL/m06300.cl | 8 -------- OpenCL/m07000_a0.cl | 4 ---- OpenCL/m07000_a3.cl | 4 ---- OpenCL/m07300_a0.cl | 4 ---- OpenCL/m07300_a1.cl | 4 ---- OpenCL/m07300_a3.cl | 4 ---- OpenCL/m07400.cl | 4 ---- OpenCL/m07500_a0.cl | 4 ---- OpenCL/m07500_a3.cl | 4 ---- OpenCL/m07900.cl | 2 -- OpenCL/m08100_a0.cl | 4 ---- OpenCL/m08100_a3.cl | 4 ---- OpenCL/m08300_a0.cl | 12 ------------ OpenCL/m08300_a1.cl | 8 -------- OpenCL/m08300_a3.cl | 12 ------------ OpenCL/m08400_a0.cl | 4 ---- OpenCL/m08400_a3.cl | 4 ---- OpenCL/m08900.cl | 4 ---- OpenCL/m09900_a0.cl | 4 ---- OpenCL/m09900_a3.cl | 4 ---- OpenCL/m10700.cl | 2 -- OpenCL/m10800_a0.cl | 4 ---- OpenCL/m10800_a3.cl | 4 ---- OpenCL/m11000_a0.cl | 4 ---- OpenCL/m11000_a3.cl | 4 ---- OpenCL/m11100_a0.cl | 4 ---- OpenCL/m11100_a3.cl | 4 ---- OpenCL/m11200_a0.cl | 4 ---- OpenCL/m11200_a3.cl | 4 ---- OpenCL/m11400_a0.cl | 4 ---- OpenCL/m11400_a3.cl | 8 -------- OpenCL/m11600.cl | 2 -- OpenCL/m12400.cl | 2 -- OpenCL/m12600_a0.cl | 4 ---- OpenCL/m12600_a3.cl | 4 ---- OpenCL/m13100_a0.cl | 4 ---- OpenCL/m13100_a3.cl | 4 ---- OpenCL/m13300_a0.cl | 4 ---- OpenCL/m13300_a3.cl | 4 ---- OpenCL/m15700.cl | 4 ---- 180 files changed, 978 deletions(-) diff --git a/OpenCL/m00000_a0.cl b/OpenCL/m00000_a0.cl index a30464522..82690362f 100644 --- a/OpenCL/m00000_a0.cl +++ b/OpenCL/m00000_a0.cl @@ -39,8 +39,6 @@ __kernel void m00000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -104,8 +102,6 @@ __kernel void m00000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00000_a3.cl b/OpenCL/m00000_a3.cl index 175c759fd..cb172a09d 100644 --- a/OpenCL/m00000_a3.cl +++ b/OpenCL/m00000_a3.cl @@ -37,8 +37,6 @@ __kernel void m00000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -108,8 +106,6 @@ __kernel void m00000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00010_a0.cl b/OpenCL/m00010_a0.cl index f353def68..76e002021 100644 --- a/OpenCL/m00010_a0.cl +++ b/OpenCL/m00010_a0.cl @@ -39,8 +39,6 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -119,8 +115,6 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -132,8 +126,6 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00010_a1.cl b/OpenCL/m00010_a1.cl index 7a18b79d6..b970e65e5 100644 --- a/OpenCL/m00010_a1.cl +++ b/OpenCL/m00010_a1.cl @@ -37,8 +37,6 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -106,8 +104,6 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m00010_a3.cl b/OpenCL/m00010_a3.cl index 304bad19d..02bbb9768 100644 --- a/OpenCL/m00010_a3.cl +++ b/OpenCL/m00010_a3.cl @@ -37,8 +37,6 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -123,8 +119,6 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -136,8 +130,6 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00020_a0.cl b/OpenCL/m00020_a0.cl index ec7f7187d..4cd3e5e84 100644 --- a/OpenCL/m00020_a0.cl +++ b/OpenCL/m00020_a0.cl @@ -39,8 +39,6 @@ __kernel void m00020_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -108,8 +106,6 @@ __kernel void m00020_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m00020_a3.cl b/OpenCL/m00020_a3.cl index fb46e8374..c22d90731 100644 --- a/OpenCL/m00020_a3.cl +++ b/OpenCL/m00020_a3.cl @@ -37,8 +37,6 @@ __kernel void m00020_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -114,8 +112,6 @@ __kernel void m00020_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m00030_a0.cl b/OpenCL/m00030_a0.cl index 81fe98a62..349802d62 100644 --- a/OpenCL/m00030_a0.cl +++ b/OpenCL/m00030_a0.cl @@ -39,8 +39,6 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -119,8 +115,6 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -132,8 +126,6 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00030_a1.cl b/OpenCL/m00030_a1.cl index 2160c48e9..babf8745d 100644 --- a/OpenCL/m00030_a1.cl +++ b/OpenCL/m00030_a1.cl @@ -37,8 +37,6 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -106,8 +104,6 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m00030_a3.cl b/OpenCL/m00030_a3.cl index 10c2b0c00..d750b40cb 100644 --- a/OpenCL/m00030_a3.cl +++ b/OpenCL/m00030_a3.cl @@ -37,8 +37,6 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -123,8 +119,6 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -136,8 +130,6 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00040_a0.cl b/OpenCL/m00040_a0.cl index bf4aa3ff9..af4b80719 100644 --- a/OpenCL/m00040_a0.cl +++ b/OpenCL/m00040_a0.cl @@ -39,8 +39,6 @@ __kernel void m00040_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -108,8 +106,6 @@ __kernel void m00040_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m00040_a3.cl b/OpenCL/m00040_a3.cl index be5995613..af341b374 100644 --- a/OpenCL/m00040_a3.cl +++ b/OpenCL/m00040_a3.cl @@ -37,8 +37,6 @@ __kernel void m00040_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -114,8 +112,6 @@ __kernel void m00040_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m00050_a0.cl b/OpenCL/m00050_a0.cl index 263ef488f..51e5e14f6 100644 --- a/OpenCL/m00050_a0.cl +++ b/OpenCL/m00050_a0.cl @@ -39,8 +39,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -117,8 +113,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -130,8 +124,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00050_a1.cl b/OpenCL/m00050_a1.cl index de0739fb6..702ebebbe 100644 --- a/OpenCL/m00050_a1.cl +++ b/OpenCL/m00050_a1.cl @@ -37,8 +37,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -135,8 +131,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -148,8 +142,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00050_a3.cl b/OpenCL/m00050_a3.cl index 09f223a60..209ecdb74 100644 --- a/OpenCL/m00050_a3.cl +++ b/OpenCL/m00050_a3.cl @@ -37,8 +37,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -121,8 +117,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -134,8 +128,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00060_a0.cl b/OpenCL/m00060_a0.cl index ea4b96827..75369b728 100644 --- a/OpenCL/m00060_a0.cl +++ b/OpenCL/m00060_a0.cl @@ -39,8 +39,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_hmac_ctx_t ctx0; @@ -119,8 +115,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -132,8 +126,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_hmac_ctx_t ctx0; diff --git a/OpenCL/m00060_a1.cl b/OpenCL/m00060_a1.cl index a0e4f0a23..6cff7c337 100644 --- a/OpenCL/m00060_a1.cl +++ b/OpenCL/m00060_a1.cl @@ -37,8 +37,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_hmac_ctx_t ctx0; @@ -137,8 +133,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -150,8 +144,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_hmac_ctx_t ctx0; diff --git a/OpenCL/m00060_a3.cl b/OpenCL/m00060_a3.cl index 1c1d79a29..c1f165249 100644 --- a/OpenCL/m00060_a3.cl +++ b/OpenCL/m00060_a3.cl @@ -37,8 +37,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_hmac_ctx_vector_t ctx0; @@ -123,8 +119,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -136,8 +130,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_hmac_ctx_vector_t ctx0; diff --git a/OpenCL/m00100_a0.cl b/OpenCL/m00100_a0.cl index a8bd7ec57..c14eaef04 100644 --- a/OpenCL/m00100_a0.cl +++ b/OpenCL/m00100_a0.cl @@ -39,8 +39,6 @@ __kernel void m00100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -104,8 +102,6 @@ __kernel void m00100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00100_a3.cl b/OpenCL/m00100_a3.cl index 50ecd137b..f125ff993 100644 --- a/OpenCL/m00100_a3.cl +++ b/OpenCL/m00100_a3.cl @@ -37,8 +37,6 @@ __kernel void m00100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -108,8 +106,6 @@ __kernel void m00100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00110_a0.cl b/OpenCL/m00110_a0.cl index da9c38d03..030b0ca2e 100644 --- a/OpenCL/m00110_a0.cl +++ b/OpenCL/m00110_a0.cl @@ -39,8 +39,6 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -119,8 +115,6 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -132,8 +126,6 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00110_a1.cl b/OpenCL/m00110_a1.cl index 3d8987f72..e9a349591 100644 --- a/OpenCL/m00110_a1.cl +++ b/OpenCL/m00110_a1.cl @@ -37,8 +37,6 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -106,8 +104,6 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m00110_a3.cl b/OpenCL/m00110_a3.cl index 00418eafc..6320923dd 100644 --- a/OpenCL/m00110_a3.cl +++ b/OpenCL/m00110_a3.cl @@ -37,8 +37,6 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -123,8 +119,6 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -136,8 +130,6 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00120_a0.cl b/OpenCL/m00120_a0.cl index 6849a85a8..78dca9c65 100644 --- a/OpenCL/m00120_a0.cl +++ b/OpenCL/m00120_a0.cl @@ -39,8 +39,6 @@ __kernel void m00120_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -108,8 +106,6 @@ __kernel void m00120_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m00120_a3.cl b/OpenCL/m00120_a3.cl index 178dd369e..04ee961f2 100644 --- a/OpenCL/m00120_a3.cl +++ b/OpenCL/m00120_a3.cl @@ -37,8 +37,6 @@ __kernel void m00120_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -114,8 +112,6 @@ __kernel void m00120_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m00130_a0.cl b/OpenCL/m00130_a0.cl index ae3d7372c..859b3c2c1 100644 --- a/OpenCL/m00130_a0.cl +++ b/OpenCL/m00130_a0.cl @@ -39,8 +39,6 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -119,8 +115,6 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -132,8 +126,6 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00130_a1.cl b/OpenCL/m00130_a1.cl index 2f7f11bd4..4d6b8e020 100644 --- a/OpenCL/m00130_a1.cl +++ b/OpenCL/m00130_a1.cl @@ -37,8 +37,6 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -106,8 +104,6 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m00130_a3.cl b/OpenCL/m00130_a3.cl index 8024f2233..b1cf6ddcc 100644 --- a/OpenCL/m00130_a3.cl +++ b/OpenCL/m00130_a3.cl @@ -37,8 +37,6 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -124,8 +120,6 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -137,8 +131,6 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00140_a0.cl b/OpenCL/m00140_a0.cl index 1b8c5e717..aa9c8b04b 100644 --- a/OpenCL/m00140_a0.cl +++ b/OpenCL/m00140_a0.cl @@ -39,8 +39,6 @@ __kernel void m00140_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -108,8 +106,6 @@ __kernel void m00140_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m00140_a3.cl b/OpenCL/m00140_a3.cl index 75d22495b..91ad701ac 100644 --- a/OpenCL/m00140_a3.cl +++ b/OpenCL/m00140_a3.cl @@ -37,8 +37,6 @@ __kernel void m00140_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -114,8 +112,6 @@ __kernel void m00140_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m00150_a0.cl b/OpenCL/m00150_a0.cl index 957d2a50b..3e4d11f88 100644 --- a/OpenCL/m00150_a0.cl +++ b/OpenCL/m00150_a0.cl @@ -39,8 +39,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -117,8 +113,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -130,8 +124,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00150_a1.cl b/OpenCL/m00150_a1.cl index 96153a25c..9b55986b6 100644 --- a/OpenCL/m00150_a1.cl +++ b/OpenCL/m00150_a1.cl @@ -37,8 +37,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -135,8 +131,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -148,8 +142,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00150_a3.cl b/OpenCL/m00150_a3.cl index 27c7dc7dd..195218abf 100644 --- a/OpenCL/m00150_a3.cl +++ b/OpenCL/m00150_a3.cl @@ -37,8 +37,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -121,8 +117,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -134,8 +128,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00160_a0.cl b/OpenCL/m00160_a0.cl index c87a1bab7..d32cdac55 100644 --- a/OpenCL/m00160_a0.cl +++ b/OpenCL/m00160_a0.cl @@ -39,8 +39,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_hmac_ctx_t ctx0; @@ -119,8 +115,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -132,8 +126,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_hmac_ctx_t ctx0; diff --git a/OpenCL/m00160_a1.cl b/OpenCL/m00160_a1.cl index 6a4907c55..b8329a326 100644 --- a/OpenCL/m00160_a1.cl +++ b/OpenCL/m00160_a1.cl @@ -37,8 +37,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_hmac_ctx_t ctx0; @@ -137,8 +133,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -150,8 +144,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_hmac_ctx_t ctx0; diff --git a/OpenCL/m00160_a3.cl b/OpenCL/m00160_a3.cl index cdcca3d64..1712e8064 100644 --- a/OpenCL/m00160_a3.cl +++ b/OpenCL/m00160_a3.cl @@ -37,8 +37,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_hmac_ctx_vector_t ctx0; @@ -123,8 +119,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -136,8 +130,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_hmac_ctx_vector_t ctx0; diff --git a/OpenCL/m00300_a0.cl b/OpenCL/m00300_a0.cl index 04e7f4940..d05d17076 100644 --- a/OpenCL/m00300_a0.cl +++ b/OpenCL/m00300_a0.cl @@ -39,8 +39,6 @@ __kernel void m00300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -127,8 +125,6 @@ __kernel void m00300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00300_a3.cl b/OpenCL/m00300_a3.cl index 6a9763f3f..d68b2608c 100644 --- a/OpenCL/m00300_a3.cl +++ b/OpenCL/m00300_a3.cl @@ -37,8 +37,6 @@ __kernel void m00300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -131,8 +129,6 @@ __kernel void m00300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00400.cl b/OpenCL/m00400.cl index 821e4a02a..5d3ad84b4 100644 --- a/OpenCL/m00400.cl +++ b/OpenCL/m00400.cl @@ -76,8 +76,6 @@ __kernel void m00400_loop (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } u32 digest[4]; diff --git a/OpenCL/m00500.cl b/OpenCL/m00500.cl index 54422ec96..80371e61d 100644 --- a/OpenCL/m00500.cl +++ b/OpenCL/m00500.cl @@ -40,8 +40,6 @@ __kernel void m00500_init (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -53,8 +51,6 @@ __kernel void m00500_init (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -154,8 +150,6 @@ __kernel void m00500_loop (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -167,8 +161,6 @@ __kernel void m00500_loop (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00900_a0.cl b/OpenCL/m00900_a0.cl index cf19e1c05..a58b5f934 100644 --- a/OpenCL/m00900_a0.cl +++ b/OpenCL/m00900_a0.cl @@ -39,8 +39,6 @@ __kernel void m00900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -104,8 +102,6 @@ __kernel void m00900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m00900_a3.cl b/OpenCL/m00900_a3.cl index 480a1e3c4..52ab4da60 100644 --- a/OpenCL/m00900_a3.cl +++ b/OpenCL/m00900_a3.cl @@ -37,8 +37,6 @@ __kernel void m00900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -108,8 +106,6 @@ __kernel void m00900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01000_a0.cl b/OpenCL/m01000_a0.cl index d01ddc0d5..c28b005e6 100644 --- a/OpenCL/m01000_a0.cl +++ b/OpenCL/m01000_a0.cl @@ -39,8 +39,6 @@ __kernel void m01000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -104,8 +102,6 @@ __kernel void m01000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01000_a3.cl b/OpenCL/m01000_a3.cl index a9a421686..d4048fffb 100644 --- a/OpenCL/m01000_a3.cl +++ b/OpenCL/m01000_a3.cl @@ -37,8 +37,6 @@ __kernel void m01000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -108,8 +106,6 @@ __kernel void m01000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01100_a0.cl b/OpenCL/m01100_a0.cl index 085e29e18..9dc3d1638 100644 --- a/OpenCL/m01100_a0.cl +++ b/OpenCL/m01100_a0.cl @@ -39,8 +39,6 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -132,8 +128,6 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -145,8 +139,6 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01100_a1.cl b/OpenCL/m01100_a1.cl index 6c7ef2d9e..893122399 100644 --- a/OpenCL/m01100_a1.cl +++ b/OpenCL/m01100_a1.cl @@ -37,8 +37,6 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md4_ctx_t ctx0; @@ -119,8 +117,6 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md4_ctx_t ctx0; diff --git a/OpenCL/m01100_a3.cl b/OpenCL/m01100_a3.cl index 461b834a8..03e59c812 100644 --- a/OpenCL/m01100_a3.cl +++ b/OpenCL/m01100_a3.cl @@ -37,8 +37,6 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -136,8 +132,6 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -149,8 +143,6 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01300_a0.cl b/OpenCL/m01300_a0.cl index 7607c95f8..588e58786 100644 --- a/OpenCL/m01300_a0.cl +++ b/OpenCL/m01300_a0.cl @@ -39,8 +39,6 @@ __kernel void m01300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -104,8 +102,6 @@ __kernel void m01300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01300_a3.cl b/OpenCL/m01300_a3.cl index cbb1203a8..d756fd264 100644 --- a/OpenCL/m01300_a3.cl +++ b/OpenCL/m01300_a3.cl @@ -37,8 +37,6 @@ __kernel void m01300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -108,8 +106,6 @@ __kernel void m01300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01400_a0.cl b/OpenCL/m01400_a0.cl index 3f07773af..b26194a63 100644 --- a/OpenCL/m01400_a0.cl +++ b/OpenCL/m01400_a0.cl @@ -39,8 +39,6 @@ __kernel void m01400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -104,8 +102,6 @@ __kernel void m01400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01400_a3.cl b/OpenCL/m01400_a3.cl index 6595c7b9a..3cba71a25 100644 --- a/OpenCL/m01400_a3.cl +++ b/OpenCL/m01400_a3.cl @@ -37,8 +37,6 @@ __kernel void m01400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -108,8 +106,6 @@ __kernel void m01400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01410_a0.cl b/OpenCL/m01410_a0.cl index 185c9e132..b517e02bd 100644 --- a/OpenCL/m01410_a0.cl +++ b/OpenCL/m01410_a0.cl @@ -39,8 +39,6 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -119,8 +115,6 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -132,8 +126,6 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01410_a1.cl b/OpenCL/m01410_a1.cl index bf7a01885..b7d22426c 100644 --- a/OpenCL/m01410_a1.cl +++ b/OpenCL/m01410_a1.cl @@ -37,8 +37,6 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_ctx_t ctx0; @@ -106,8 +104,6 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_ctx_t ctx0; diff --git a/OpenCL/m01410_a3.cl b/OpenCL/m01410_a3.cl index a1a7bd150..dd860cce1 100644 --- a/OpenCL/m01410_a3.cl +++ b/OpenCL/m01410_a3.cl @@ -37,8 +37,6 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -123,8 +119,6 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -136,8 +130,6 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01420_a0.cl b/OpenCL/m01420_a0.cl index 4733245f6..6adc7367f 100644 --- a/OpenCL/m01420_a0.cl +++ b/OpenCL/m01420_a0.cl @@ -39,8 +39,6 @@ __kernel void m01420_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_ctx_t ctx0; @@ -108,8 +106,6 @@ __kernel void m01420_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_ctx_t ctx0; diff --git a/OpenCL/m01420_a3.cl b/OpenCL/m01420_a3.cl index 02784ecd0..11526045b 100644 --- a/OpenCL/m01420_a3.cl +++ b/OpenCL/m01420_a3.cl @@ -37,8 +37,6 @@ __kernel void m01420_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_ctx_t ctx0; @@ -114,8 +112,6 @@ __kernel void m01420_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_ctx_t ctx0; diff --git a/OpenCL/m01430_a0.cl b/OpenCL/m01430_a0.cl index 45ef46ba9..e3236eef1 100644 --- a/OpenCL/m01430_a0.cl +++ b/OpenCL/m01430_a0.cl @@ -39,8 +39,6 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -119,8 +115,6 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -132,8 +126,6 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01430_a1.cl b/OpenCL/m01430_a1.cl index 2d0efa682..8eef1f177 100644 --- a/OpenCL/m01430_a1.cl +++ b/OpenCL/m01430_a1.cl @@ -37,8 +37,6 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_ctx_t ctx0; @@ -106,8 +104,6 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_ctx_t ctx0; diff --git a/OpenCL/m01430_a3.cl b/OpenCL/m01430_a3.cl index 2a6ccc3f8..3259e7c17 100644 --- a/OpenCL/m01430_a3.cl +++ b/OpenCL/m01430_a3.cl @@ -37,8 +37,6 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -124,8 +120,6 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -137,8 +131,6 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01440_a0.cl b/OpenCL/m01440_a0.cl index 4dcf235e0..271d4367e 100644 --- a/OpenCL/m01440_a0.cl +++ b/OpenCL/m01440_a0.cl @@ -39,8 +39,6 @@ __kernel void m01440_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_ctx_t ctx0; @@ -108,8 +106,6 @@ __kernel void m01440_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_ctx_t ctx0; diff --git a/OpenCL/m01440_a3.cl b/OpenCL/m01440_a3.cl index a01604ac3..316e0d156 100644 --- a/OpenCL/m01440_a3.cl +++ b/OpenCL/m01440_a3.cl @@ -37,8 +37,6 @@ __kernel void m01440_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_ctx_t ctx0; @@ -114,8 +112,6 @@ __kernel void m01440_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_ctx_t ctx0; diff --git a/OpenCL/m01450_a0.cl b/OpenCL/m01450_a0.cl index 21fa01b0c..e1462f303 100644 --- a/OpenCL/m01450_a0.cl +++ b/OpenCL/m01450_a0.cl @@ -39,8 +39,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -117,8 +113,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -130,8 +124,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01450_a1.cl b/OpenCL/m01450_a1.cl index 25e81d266..7f7c9f635 100644 --- a/OpenCL/m01450_a1.cl +++ b/OpenCL/m01450_a1.cl @@ -37,8 +37,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -135,8 +131,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -148,8 +142,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01450_a3.cl b/OpenCL/m01450_a3.cl index c55b53eae..35bd59a2a 100644 --- a/OpenCL/m01450_a3.cl +++ b/OpenCL/m01450_a3.cl @@ -37,8 +37,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -121,8 +117,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -134,8 +128,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01460_a0.cl b/OpenCL/m01460_a0.cl index 0b14e52e9..d51c28cb6 100644 --- a/OpenCL/m01460_a0.cl +++ b/OpenCL/m01460_a0.cl @@ -39,8 +39,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_hmac_ctx_t ctx0; @@ -119,8 +115,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -132,8 +126,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_hmac_ctx_t ctx0; diff --git a/OpenCL/m01460_a1.cl b/OpenCL/m01460_a1.cl index d4b0fdb8b..61c30c095 100644 --- a/OpenCL/m01460_a1.cl +++ b/OpenCL/m01460_a1.cl @@ -37,8 +37,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_hmac_ctx_t ctx0; @@ -137,8 +133,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -150,8 +144,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_hmac_ctx_t ctx0; diff --git a/OpenCL/m01460_a3.cl b/OpenCL/m01460_a3.cl index 646ed4732..d7ac9bedf 100644 --- a/OpenCL/m01460_a3.cl +++ b/OpenCL/m01460_a3.cl @@ -37,8 +37,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_hmac_ctx_vector_t ctx0; @@ -123,8 +119,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -136,8 +130,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha256_hmac_ctx_vector_t ctx0; diff --git a/OpenCL/m01600.cl b/OpenCL/m01600.cl index dc110a813..bdc7e5334 100644 --- a/OpenCL/m01600.cl +++ b/OpenCL/m01600.cl @@ -41,8 +41,6 @@ __kernel void m01600_init (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -54,8 +52,6 @@ __kernel void m01600_init (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -156,8 +152,6 @@ __kernel void m01600_loop (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -169,8 +163,6 @@ __kernel void m01600_loop (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01700_a0.cl b/OpenCL/m01700_a0.cl index 01649a999..72e3f0f2b 100644 --- a/OpenCL/m01700_a0.cl +++ b/OpenCL/m01700_a0.cl @@ -39,8 +39,6 @@ __kernel void m01700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -104,8 +102,6 @@ __kernel void m01700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01700_a3.cl b/OpenCL/m01700_a3.cl index e35083c8b..80f66afb2 100644 --- a/OpenCL/m01700_a3.cl +++ b/OpenCL/m01700_a3.cl @@ -37,8 +37,6 @@ __kernel void m01700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -108,8 +106,6 @@ __kernel void m01700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01710_a0.cl b/OpenCL/m01710_a0.cl index 593be6b7d..acd32a0f3 100644 --- a/OpenCL/m01710_a0.cl +++ b/OpenCL/m01710_a0.cl @@ -39,8 +39,6 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -119,8 +115,6 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -132,8 +126,6 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01710_a1.cl b/OpenCL/m01710_a1.cl index 485adbb32..b09354e90 100644 --- a/OpenCL/m01710_a1.cl +++ b/OpenCL/m01710_a1.cl @@ -37,8 +37,6 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_ctx_t ctx0; @@ -106,8 +104,6 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_ctx_t ctx0; diff --git a/OpenCL/m01710_a3.cl b/OpenCL/m01710_a3.cl index 6da78255b..f54ce7355 100644 --- a/OpenCL/m01710_a3.cl +++ b/OpenCL/m01710_a3.cl @@ -37,8 +37,6 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -123,8 +119,6 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -136,8 +130,6 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01720_a0.cl b/OpenCL/m01720_a0.cl index adcf3501b..e80c9426c 100644 --- a/OpenCL/m01720_a0.cl +++ b/OpenCL/m01720_a0.cl @@ -39,8 +39,6 @@ __kernel void m01720_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_ctx_t ctx0; @@ -108,8 +106,6 @@ __kernel void m01720_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_ctx_t ctx0; diff --git a/OpenCL/m01720_a3.cl b/OpenCL/m01720_a3.cl index e6bb100d3..b28688280 100644 --- a/OpenCL/m01720_a3.cl +++ b/OpenCL/m01720_a3.cl @@ -37,8 +37,6 @@ __kernel void m01720_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_ctx_t ctx0; @@ -114,8 +112,6 @@ __kernel void m01720_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_ctx_t ctx0; diff --git a/OpenCL/m01730_a0.cl b/OpenCL/m01730_a0.cl index f4a8951a8..16a6753a7 100644 --- a/OpenCL/m01730_a0.cl +++ b/OpenCL/m01730_a0.cl @@ -39,8 +39,6 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -119,8 +115,6 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -132,8 +126,6 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01730_a1.cl b/OpenCL/m01730_a1.cl index 7529dd71d..a937af113 100644 --- a/OpenCL/m01730_a1.cl +++ b/OpenCL/m01730_a1.cl @@ -37,8 +37,6 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_ctx_t ctx0; @@ -106,8 +104,6 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_ctx_t ctx0; diff --git a/OpenCL/m01730_a3.cl b/OpenCL/m01730_a3.cl index c408f8105..90d53fd56 100644 --- a/OpenCL/m01730_a3.cl +++ b/OpenCL/m01730_a3.cl @@ -37,8 +37,6 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -124,8 +120,6 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -137,8 +131,6 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01740_a0.cl b/OpenCL/m01740_a0.cl index d489d6e67..0c76ab222 100644 --- a/OpenCL/m01740_a0.cl +++ b/OpenCL/m01740_a0.cl @@ -39,8 +39,6 @@ __kernel void m01740_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_ctx_t ctx0; @@ -108,8 +106,6 @@ __kernel void m01740_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_ctx_t ctx0; diff --git a/OpenCL/m01740_a3.cl b/OpenCL/m01740_a3.cl index 6ae09ef93..0e879be01 100644 --- a/OpenCL/m01740_a3.cl +++ b/OpenCL/m01740_a3.cl @@ -37,8 +37,6 @@ __kernel void m01740_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_ctx_t ctx0; @@ -114,8 +112,6 @@ __kernel void m01740_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_ctx_t ctx0; diff --git a/OpenCL/m01750_a0.cl b/OpenCL/m01750_a0.cl index 0ff9c0346..b3e12bff9 100644 --- a/OpenCL/m01750_a0.cl +++ b/OpenCL/m01750_a0.cl @@ -39,8 +39,6 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -117,8 +113,6 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -130,8 +124,6 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01750_a1.cl b/OpenCL/m01750_a1.cl index d6d9da0fd..430d9a7ce 100644 --- a/OpenCL/m01750_a1.cl +++ b/OpenCL/m01750_a1.cl @@ -37,8 +37,6 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -135,8 +131,6 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -148,8 +142,6 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01750_a3.cl b/OpenCL/m01750_a3.cl index 28c6a4b38..a9c59d030 100644 --- a/OpenCL/m01750_a3.cl +++ b/OpenCL/m01750_a3.cl @@ -37,8 +37,6 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -121,8 +117,6 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -134,8 +128,6 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m01760_a0.cl b/OpenCL/m01760_a0.cl index f977db7ac..460e1cb3f 100644 --- a/OpenCL/m01760_a0.cl +++ b/OpenCL/m01760_a0.cl @@ -39,8 +39,6 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_hmac_ctx_t ctx0; @@ -119,8 +115,6 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -132,8 +126,6 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_hmac_ctx_t ctx0; diff --git a/OpenCL/m01760_a1.cl b/OpenCL/m01760_a1.cl index 5403bcc5d..1a558b84f 100644 --- a/OpenCL/m01760_a1.cl +++ b/OpenCL/m01760_a1.cl @@ -37,8 +37,6 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_hmac_ctx_t ctx0; @@ -137,8 +133,6 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -150,8 +144,6 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_hmac_ctx_t ctx0; diff --git a/OpenCL/m01760_a3.cl b/OpenCL/m01760_a3.cl index f3b80a484..4bb4c0d78 100644 --- a/OpenCL/m01760_a3.cl +++ b/OpenCL/m01760_a3.cl @@ -37,8 +37,6 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_hmac_ctx_vector_t ctx0; @@ -123,8 +119,6 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -136,8 +130,6 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha512_hmac_ctx_vector_t ctx0; diff --git a/OpenCL/m01800.cl b/OpenCL/m01800.cl index 46270e3dc..02360d45e 100644 --- a/OpenCL/m01800.cl +++ b/OpenCL/m01800.cl @@ -36,8 +36,6 @@ __kernel void m01800_init (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } for (int idx = 0; idx < pw_lenv; idx++) @@ -54,8 +52,6 @@ __kernel void m01800_init (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } for (int idx = 0; idx < salt_lenv; idx++) diff --git a/OpenCL/m02610_a0.cl b/OpenCL/m02610_a0.cl index a09b5b41e..d86d026c8 100644 --- a/OpenCL/m02610_a0.cl +++ b/OpenCL/m02610_a0.cl @@ -69,8 +69,6 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -82,8 +80,6 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -197,8 +193,6 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -210,8 +204,6 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m02610_a1.cl b/OpenCL/m02610_a1.cl index 3fbdd72f9..0cfd702fa 100644 --- a/OpenCL/m02610_a1.cl +++ b/OpenCL/m02610_a1.cl @@ -67,8 +67,6 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -184,8 +182,6 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m02610_a3.cl b/OpenCL/m02610_a3.cl index 442dfe070..441ec66fd 100644 --- a/OpenCL/m02610_a3.cl +++ b/OpenCL/m02610_a3.cl @@ -67,8 +67,6 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -80,8 +78,6 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -201,8 +197,6 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -214,8 +208,6 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m02810_a0.cl b/OpenCL/m02810_a0.cl index 502105952..c1e047201 100644 --- a/OpenCL/m02810_a0.cl +++ b/OpenCL/m02810_a0.cl @@ -69,8 +69,6 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = 32; @@ -82,8 +80,6 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -220,8 +216,6 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = 32; @@ -233,8 +227,6 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m02810_a1.cl b/OpenCL/m02810_a1.cl index df010bf8d..c14141814 100644 --- a/OpenCL/m02810_a1.cl +++ b/OpenCL/m02810_a1.cl @@ -67,8 +67,6 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -207,8 +205,6 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m02810_a3.cl b/OpenCL/m02810_a3.cl index f46abd13c..5cd8bbc94 100644 --- a/OpenCL/m02810_a3.cl +++ b/OpenCL/m02810_a3.cl @@ -67,8 +67,6 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = 32; @@ -80,8 +78,6 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -224,8 +220,6 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = 32; @@ -237,8 +231,6 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m03710_a0.cl b/OpenCL/m03710_a0.cl index 89b894741..1401ad716 100644 --- a/OpenCL/m03710_a0.cl +++ b/OpenCL/m03710_a0.cl @@ -69,8 +69,6 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -82,8 +80,6 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -210,8 +206,6 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -223,8 +217,6 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m03710_a1.cl b/OpenCL/m03710_a1.cl index 16e7caa5e..18ed573ad 100644 --- a/OpenCL/m03710_a1.cl +++ b/OpenCL/m03710_a1.cl @@ -67,8 +67,6 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -197,8 +195,6 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m03710_a3.cl b/OpenCL/m03710_a3.cl index 8c35c9263..e17da66f9 100644 --- a/OpenCL/m03710_a3.cl +++ b/OpenCL/m03710_a3.cl @@ -67,8 +67,6 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -80,8 +78,6 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -214,8 +210,6 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -227,8 +221,6 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m03800_a0.cl b/OpenCL/m03800_a0.cl index 42344114b..8608dc55c 100644 --- a/OpenCL/m03800_a0.cl +++ b/OpenCL/m03800_a0.cl @@ -39,8 +39,6 @@ __kernel void m03800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m03800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -123,8 +119,6 @@ __kernel void m03800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -136,8 +130,6 @@ __kernel void m03800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m03800_a1.cl b/OpenCL/m03800_a1.cl index c594eee3d..31007631a 100644 --- a/OpenCL/m03800_a1.cl +++ b/OpenCL/m03800_a1.cl @@ -37,8 +37,6 @@ __kernel void m03800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -108,8 +106,6 @@ __kernel void m03800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m03800_a3.cl b/OpenCL/m03800_a3.cl index 00275dbeb..5736596bc 100644 --- a/OpenCL/m03800_a3.cl +++ b/OpenCL/m03800_a3.cl @@ -37,8 +37,6 @@ __kernel void m03800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m03800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -129,8 +125,6 @@ __kernel void m03800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -142,8 +136,6 @@ __kernel void m03800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m03910_a0.cl b/OpenCL/m03910_a0.cl index 1139e0cb6..a296705c5 100644 --- a/OpenCL/m03910_a0.cl +++ b/OpenCL/m03910_a0.cl @@ -69,8 +69,6 @@ __kernel void m03910_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = 32; @@ -82,8 +80,6 @@ __kernel void m03910_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -220,8 +216,6 @@ __kernel void m03910_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = 32; @@ -233,8 +227,6 @@ __kernel void m03910_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m03910_a1.cl b/OpenCL/m03910_a1.cl index 40c9071dc..1a082b9df 100644 --- a/OpenCL/m03910_a1.cl +++ b/OpenCL/m03910_a1.cl @@ -67,8 +67,6 @@ __kernel void m03910_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -207,8 +205,6 @@ __kernel void m03910_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m03910_a3.cl b/OpenCL/m03910_a3.cl index 48e7ae215..95ca4a3f3 100644 --- a/OpenCL/m03910_a3.cl +++ b/OpenCL/m03910_a3.cl @@ -67,8 +67,6 @@ __kernel void m03910_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = 32; @@ -80,8 +78,6 @@ __kernel void m03910_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -224,8 +220,6 @@ __kernel void m03910_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = 32; @@ -237,8 +231,6 @@ __kernel void m03910_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m04010_a0.cl b/OpenCL/m04010_a0.cl index 454997290..e1365e632 100644 --- a/OpenCL/m04010_a0.cl +++ b/OpenCL/m04010_a0.cl @@ -69,8 +69,6 @@ __kernel void m04010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -197,8 +195,6 @@ __kernel void m04010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m04010_a3.cl b/OpenCL/m04010_a3.cl index 4822b707b..91b54e48f 100644 --- a/OpenCL/m04010_a3.cl +++ b/OpenCL/m04010_a3.cl @@ -67,8 +67,6 @@ __kernel void m04010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -205,8 +203,6 @@ __kernel void m04010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m04110_a0.cl b/OpenCL/m04110_a0.cl index 4bcb4def6..ed85497a3 100644 --- a/OpenCL/m04110_a0.cl +++ b/OpenCL/m04110_a0.cl @@ -69,8 +69,6 @@ __kernel void m04110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -82,8 +80,6 @@ __kernel void m04110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -214,8 +210,6 @@ __kernel void m04110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -227,8 +221,6 @@ __kernel void m04110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m04110_a1.cl b/OpenCL/m04110_a1.cl index 57684942a..5bce915db 100644 --- a/OpenCL/m04110_a1.cl +++ b/OpenCL/m04110_a1.cl @@ -67,8 +67,6 @@ __kernel void m04110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -201,8 +199,6 @@ __kernel void m04110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m04110_a3.cl b/OpenCL/m04110_a3.cl index f466a3178..57ec737f4 100644 --- a/OpenCL/m04110_a3.cl +++ b/OpenCL/m04110_a3.cl @@ -67,8 +67,6 @@ __kernel void m04110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -80,8 +78,6 @@ __kernel void m04110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -220,8 +216,6 @@ __kernel void m04110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -233,8 +227,6 @@ __kernel void m04110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m04310_a0.cl b/OpenCL/m04310_a0.cl index 31d1af318..6e2ac764e 100644 --- a/OpenCL/m04310_a0.cl +++ b/OpenCL/m04310_a0.cl @@ -69,8 +69,6 @@ __kernel void m04310_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -82,8 +80,6 @@ __kernel void m04310_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -197,8 +193,6 @@ __kernel void m04310_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -210,8 +204,6 @@ __kernel void m04310_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m04310_a1.cl b/OpenCL/m04310_a1.cl index bf33a9155..3ed91424c 100644 --- a/OpenCL/m04310_a1.cl +++ b/OpenCL/m04310_a1.cl @@ -67,8 +67,6 @@ __kernel void m04310_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -184,8 +182,6 @@ __kernel void m04310_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m04310_a3.cl b/OpenCL/m04310_a3.cl index 2f491a2fb..e265cf4a1 100644 --- a/OpenCL/m04310_a3.cl +++ b/OpenCL/m04310_a3.cl @@ -67,8 +67,6 @@ __kernel void m04310_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -80,8 +78,6 @@ __kernel void m04310_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -201,8 +197,6 @@ __kernel void m04310_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -214,8 +208,6 @@ __kernel void m04310_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m04400_a0.cl b/OpenCL/m04400_a0.cl index 877fee359..7d583393c 100644 --- a/OpenCL/m04400_a0.cl +++ b/OpenCL/m04400_a0.cl @@ -70,8 +70,6 @@ __kernel void m04400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -188,8 +186,6 @@ __kernel void m04400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m04400_a3.cl b/OpenCL/m04400_a3.cl index 598de7915..8439a874d 100644 --- a/OpenCL/m04400_a3.cl +++ b/OpenCL/m04400_a3.cl @@ -68,8 +68,6 @@ __kernel void m04400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -192,8 +190,6 @@ __kernel void m04400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m04500_a0.cl b/OpenCL/m04500_a0.cl index 5f2d0d62e..35d111c0e 100644 --- a/OpenCL/m04500_a0.cl +++ b/OpenCL/m04500_a0.cl @@ -69,8 +69,6 @@ __kernel void m04500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -187,8 +185,6 @@ __kernel void m04500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m04500_a3.cl b/OpenCL/m04500_a3.cl index c14543397..cea19b72f 100644 --- a/OpenCL/m04500_a3.cl +++ b/OpenCL/m04500_a3.cl @@ -67,8 +67,6 @@ __kernel void m04500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -191,8 +189,6 @@ __kernel void m04500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m04520_a0.cl b/OpenCL/m04520_a0.cl index 5ac22b8ac..bf8924407 100644 --- a/OpenCL/m04520_a0.cl +++ b/OpenCL/m04520_a0.cl @@ -69,8 +69,6 @@ __kernel void m04520_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -202,8 +200,6 @@ __kernel void m04520_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m04520_a3.cl b/OpenCL/m04520_a3.cl index dd633dda1..ff9a71e82 100644 --- a/OpenCL/m04520_a3.cl +++ b/OpenCL/m04520_a3.cl @@ -67,8 +67,6 @@ __kernel void m04520_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -208,8 +206,6 @@ __kernel void m04520_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m04700_a0.cl b/OpenCL/m04700_a0.cl index bb2374eaf..611834f5d 100644 --- a/OpenCL/m04700_a0.cl +++ b/OpenCL/m04700_a0.cl @@ -70,8 +70,6 @@ __kernel void m04700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -184,8 +182,6 @@ __kernel void m04700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m04700_a3.cl b/OpenCL/m04700_a3.cl index 9b00733ed..9211ec451 100644 --- a/OpenCL/m04700_a3.cl +++ b/OpenCL/m04700_a3.cl @@ -68,8 +68,6 @@ __kernel void m04700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -187,8 +185,6 @@ __kernel void m04700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m04800_a0.cl b/OpenCL/m04800_a0.cl index cd7c9cc07..fea15ad37 100644 --- a/OpenCL/m04800_a0.cl +++ b/OpenCL/m04800_a0.cl @@ -39,8 +39,6 @@ __kernel void m04800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len - 1; @@ -121,8 +119,6 @@ __kernel void m04800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len - 1; diff --git a/OpenCL/m04800_a3.cl b/OpenCL/m04800_a3.cl index e69628fbb..ed25c642c 100644 --- a/OpenCL/m04800_a3.cl +++ b/OpenCL/m04800_a3.cl @@ -37,8 +37,6 @@ __kernel void m04800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len - 1; @@ -127,8 +125,6 @@ __kernel void m04800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len - 1; diff --git a/OpenCL/m04900_a0.cl b/OpenCL/m04900_a0.cl index b323e6e6e..b1332a3de 100644 --- a/OpenCL/m04900_a0.cl +++ b/OpenCL/m04900_a0.cl @@ -39,8 +39,6 @@ __kernel void m04900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m04900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -123,8 +119,6 @@ __kernel void m04900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -136,8 +130,6 @@ __kernel void m04900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m04900_a1.cl b/OpenCL/m04900_a1.cl index a13057b0e..1a05735e3 100644 --- a/OpenCL/m04900_a1.cl +++ b/OpenCL/m04900_a1.cl @@ -37,8 +37,6 @@ __kernel void m04900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -108,8 +106,6 @@ __kernel void m04900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m04900_a3.cl b/OpenCL/m04900_a3.cl index 6b365a6f8..3f9668d61 100644 --- a/OpenCL/m04900_a3.cl +++ b/OpenCL/m04900_a3.cl @@ -37,8 +37,6 @@ __kernel void m04900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m04900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32 (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -129,8 +125,6 @@ __kernel void m04900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -142,8 +136,6 @@ __kernel void m04900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32 (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m05100_a0.cl b/OpenCL/m05100_a0.cl index 9da9df5b7..bbaf71b68 100644 --- a/OpenCL/m05100_a0.cl +++ b/OpenCL/m05100_a0.cl @@ -39,8 +39,6 @@ __kernel void m05100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -108,8 +106,6 @@ __kernel void m05100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m05100_a3.cl b/OpenCL/m05100_a3.cl index e8a25bdbc..3fcb567c9 100644 --- a/OpenCL/m05100_a3.cl +++ b/OpenCL/m05100_a3.cl @@ -37,8 +37,6 @@ __kernel void m05100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -112,8 +110,6 @@ __kernel void m05100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m05300_a0.cl b/OpenCL/m05300_a0.cl index 6c17f8865..552785041 100644 --- a/OpenCL/m05300_a0.cl +++ b/OpenCL/m05300_a0.cl @@ -39,8 +39,6 @@ __kernel void m05300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -134,8 +132,6 @@ __kernel void m05300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m05300_a1.cl b/OpenCL/m05300_a1.cl index aedad50e9..e96b61ba2 100644 --- a/OpenCL/m05300_a1.cl +++ b/OpenCL/m05300_a1.cl @@ -37,8 +37,6 @@ __kernel void m05300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -152,8 +150,6 @@ __kernel void m05300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m05300_a3.cl b/OpenCL/m05300_a3.cl index e42db4387..80891631c 100644 --- a/OpenCL/m05300_a3.cl +++ b/OpenCL/m05300_a3.cl @@ -37,8 +37,6 @@ __kernel void m05300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -138,8 +136,6 @@ __kernel void m05300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m05400_a0.cl b/OpenCL/m05400_a0.cl index 97d759c03..252e18f61 100644 --- a/OpenCL/m05400_a0.cl +++ b/OpenCL/m05400_a0.cl @@ -39,8 +39,6 @@ __kernel void m05400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -134,8 +132,6 @@ __kernel void m05400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m05400_a1.cl b/OpenCL/m05400_a1.cl index 74bd00d06..8aa75bef2 100644 --- a/OpenCL/m05400_a1.cl +++ b/OpenCL/m05400_a1.cl @@ -37,8 +37,6 @@ __kernel void m05400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -152,8 +150,6 @@ __kernel void m05400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m05400_a3.cl b/OpenCL/m05400_a3.cl index 92c014049..e40fd0cbd 100644 --- a/OpenCL/m05400_a3.cl +++ b/OpenCL/m05400_a3.cl @@ -37,8 +37,6 @@ __kernel void m05400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -138,8 +136,6 @@ __kernel void m05400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m05500_a0.cl b/OpenCL/m05500_a0.cl index 3780e2e1b..d10338f6c 100644 --- a/OpenCL/m05500_a0.cl +++ b/OpenCL/m05500_a0.cl @@ -553,8 +553,6 @@ __kernel void m05500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -698,8 +696,6 @@ __kernel void m05500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m05500_a3.cl b/OpenCL/m05500_a3.cl index 261e49457..192e3c2a7 100644 --- a/OpenCL/m05500_a3.cl +++ b/OpenCL/m05500_a3.cl @@ -550,8 +550,6 @@ __kernel void m05500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -701,8 +699,6 @@ __kernel void m05500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m05600_a0.cl b/OpenCL/m05600_a0.cl index d839d314d..1c2ac82bf 100644 --- a/OpenCL/m05600_a0.cl +++ b/OpenCL/m05600_a0.cl @@ -40,8 +40,6 @@ __kernel void m05600_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -160,8 +158,6 @@ __kernel void m05600_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m05600_a3.cl b/OpenCL/m05600_a3.cl index c963a3d66..7ec0e36a9 100644 --- a/OpenCL/m05600_a3.cl +++ b/OpenCL/m05600_a3.cl @@ -38,8 +38,6 @@ __kernel void m05600_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -164,8 +162,6 @@ __kernel void m05600_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m05800.cl b/OpenCL/m05800.cl index a447bf79e..4569e1c41 100644 --- a/OpenCL/m05800.cl +++ b/OpenCL/m05800.cl @@ -2267,8 +2267,6 @@ __kernel void m05800_loop (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32 (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -2280,8 +2278,6 @@ __kernel void m05800_loop (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32 (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } u32 digest[5]; diff --git a/OpenCL/m06000_a0.cl b/OpenCL/m06000_a0.cl index 6f4c1d1a0..cc67d8593 100644 --- a/OpenCL/m06000_a0.cl +++ b/OpenCL/m06000_a0.cl @@ -39,8 +39,6 @@ __kernel void m06000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -104,8 +102,6 @@ __kernel void m06000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m06000_a3.cl b/OpenCL/m06000_a3.cl index e3b171d63..62f84f02a 100644 --- a/OpenCL/m06000_a3.cl +++ b/OpenCL/m06000_a3.cl @@ -37,8 +37,6 @@ __kernel void m06000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -108,8 +106,6 @@ __kernel void m06000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m06100_a0.cl b/OpenCL/m06100_a0.cl index 48fcad611..5b798d0c5 100644 --- a/OpenCL/m06100_a0.cl +++ b/OpenCL/m06100_a0.cl @@ -70,8 +70,6 @@ __kernel void m06100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -166,8 +164,6 @@ __kernel void m06100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m06100_a3.cl b/OpenCL/m06100_a3.cl index da99bac26..a9bd2a358 100644 --- a/OpenCL/m06100_a3.cl +++ b/OpenCL/m06100_a3.cl @@ -68,8 +68,6 @@ __kernel void m06100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -170,8 +168,6 @@ __kernel void m06100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m06300.cl b/OpenCL/m06300.cl index dbe156e8d..0397d0362 100644 --- a/OpenCL/m06300.cl +++ b/OpenCL/m06300.cl @@ -38,8 +38,6 @@ __kernel void m06300_init (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -51,8 +49,6 @@ __kernel void m06300_init (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -146,8 +142,6 @@ __kernel void m06300_loop (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -159,8 +153,6 @@ __kernel void m06300_loop (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m07000_a0.cl b/OpenCL/m07000_a0.cl index 40201efbf..412a929ef 100644 --- a/OpenCL/m07000_a0.cl +++ b/OpenCL/m07000_a0.cl @@ -39,8 +39,6 @@ __kernel void m07000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -137,8 +135,6 @@ __kernel void m07000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m07000_a3.cl b/OpenCL/m07000_a3.cl index e80136832..6f5bf96ed 100644 --- a/OpenCL/m07000_a3.cl +++ b/OpenCL/m07000_a3.cl @@ -37,8 +37,6 @@ __kernel void m07000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -142,8 +140,6 @@ __kernel void m07000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m07300_a0.cl b/OpenCL/m07300_a0.cl index 1b78ea340..596db3a54 100644 --- a/OpenCL/m07300_a0.cl +++ b/OpenCL/m07300_a0.cl @@ -39,8 +39,6 @@ __kernel void m07300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -104,8 +102,6 @@ __kernel void m07300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m07300_a1.cl b/OpenCL/m07300_a1.cl index 3e8e63a71..775d68a0f 100644 --- a/OpenCL/m07300_a1.cl +++ b/OpenCL/m07300_a1.cl @@ -37,8 +37,6 @@ __kernel void m07300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -122,8 +120,6 @@ __kernel void m07300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m07300_a3.cl b/OpenCL/m07300_a3.cl index ff2f6f0b1..4f5532822 100644 --- a/OpenCL/m07300_a3.cl +++ b/OpenCL/m07300_a3.cl @@ -37,8 +37,6 @@ __kernel void m07300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -108,8 +106,6 @@ __kernel void m07300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m07400.cl b/OpenCL/m07400.cl index 936add9f0..9eda4d84a 100644 --- a/OpenCL/m07400.cl +++ b/OpenCL/m07400.cl @@ -36,8 +36,6 @@ __kernel void m07400_init (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } for (int idx = 0; idx < pw_lenv; idx++) @@ -54,8 +52,6 @@ __kernel void m07400_init (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = salt_bufs[salt_pos].salt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } for (int idx = 0; idx < salt_lenv; idx++) diff --git a/OpenCL/m07500_a0.cl b/OpenCL/m07500_a0.cl index e38afb3dd..72491071f 100644 --- a/OpenCL/m07500_a0.cl +++ b/OpenCL/m07500_a0.cl @@ -292,8 +292,6 @@ __kernel void m07500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } __local RC4_KEY rc4_keys[64]; @@ -370,8 +368,6 @@ __kernel void m07500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } __local RC4_KEY rc4_keys[64]; diff --git a/OpenCL/m07500_a3.cl b/OpenCL/m07500_a3.cl index 2afee5d50..9271e3924 100644 --- a/OpenCL/m07500_a3.cl +++ b/OpenCL/m07500_a3.cl @@ -302,8 +302,6 @@ __kernel void m07500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } __local RC4_KEY rc4_keys[64]; @@ -398,8 +396,6 @@ __kernel void m07500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } __local RC4_KEY rc4_keys[64]; diff --git a/OpenCL/m07900.cl b/OpenCL/m07900.cl index d50b1d4c2..eaca08e00 100644 --- a/OpenCL/m07900.cl +++ b/OpenCL/m07900.cl @@ -66,8 +66,6 @@ __kernel void m07900_loop (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } for (int idx = 0; idx < pw_lenv; idx++) diff --git a/OpenCL/m08100_a0.cl b/OpenCL/m08100_a0.cl index 8d087d790..3f3839efa 100644 --- a/OpenCL/m08100_a0.cl +++ b/OpenCL/m08100_a0.cl @@ -39,8 +39,6 @@ __kernel void m08100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -108,8 +106,6 @@ __kernel void m08100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m08100_a3.cl b/OpenCL/m08100_a3.cl index 59f0b281a..746fd240b 100644 --- a/OpenCL/m08100_a3.cl +++ b/OpenCL/m08100_a3.cl @@ -37,8 +37,6 @@ __kernel void m08100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -114,8 +112,6 @@ __kernel void m08100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m08300_a0.cl b/OpenCL/m08300_a0.cl index 39cdf7d2a..906050bef 100644 --- a/OpenCL/m08300_a0.cl +++ b/OpenCL/m08300_a0.cl @@ -39,8 +39,6 @@ __kernel void m08300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -52,8 +50,6 @@ __kernel void m08300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len_pc = salt_bufs[salt_pos].salt_len_pc; @@ -65,8 +61,6 @@ __kernel void m08300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_len_pcv; idx++) { s_pc[idx] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_iter = salt_bufs[salt_pos].salt_iter; @@ -175,8 +169,6 @@ __kernel void m08300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -188,8 +180,6 @@ __kernel void m08300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len_pc = salt_bufs[salt_pos].salt_len_pc; @@ -201,8 +191,6 @@ __kernel void m08300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_len_pcv; idx++) { s_pc[idx] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_iter = salt_bufs[salt_pos].salt_iter; diff --git a/OpenCL/m08300_a1.cl b/OpenCL/m08300_a1.cl index 57904d8d5..c5fd82a79 100644 --- a/OpenCL/m08300_a1.cl +++ b/OpenCL/m08300_a1.cl @@ -37,8 +37,6 @@ __kernel void m08300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len_pc = salt_bufs[salt_pos].salt_len_pc; @@ -50,8 +48,6 @@ __kernel void m08300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_len_pcv; idx++) { s_pc[idx] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_iter = salt_bufs[salt_pos].salt_iter; @@ -160,8 +156,6 @@ __kernel void m08300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len_pc = salt_bufs[salt_pos].salt_len_pc; @@ -173,8 +167,6 @@ __kernel void m08300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_len_pcv; idx++) { s_pc[idx] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_iter = salt_bufs[salt_pos].salt_iter; diff --git a/OpenCL/m08300_a3.cl b/OpenCL/m08300_a3.cl index 0fc7148f0..c15333c33 100644 --- a/OpenCL/m08300_a3.cl +++ b/OpenCL/m08300_a3.cl @@ -37,8 +37,6 @@ __kernel void m08300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -50,8 +48,6 @@ __kernel void m08300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32 (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len_pc = salt_bufs[salt_pos].salt_len_pc; @@ -63,8 +59,6 @@ __kernel void m08300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_len_pcv; idx++) { s_pc[idx] = swap32 (salt_bufs[salt_pos].salt_buf_pc[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_iter = salt_bufs[salt_pos].salt_iter; @@ -179,8 +173,6 @@ __kernel void m08300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -192,8 +184,6 @@ __kernel void m08300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_lenv; idx++) { s[idx] = swap32 (salt_bufs[salt_pos].salt_buf[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_len_pc = salt_bufs[salt_pos].salt_len_pc; @@ -205,8 +195,6 @@ __kernel void m08300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < salt_len_pcv; idx++) { s_pc[idx] = swap32 (salt_bufs[salt_pos].salt_buf_pc[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 salt_iter = salt_bufs[salt_pos].salt_iter; diff --git a/OpenCL/m08400_a0.cl b/OpenCL/m08400_a0.cl index 021a82f91..123887eb7 100644 --- a/OpenCL/m08400_a0.cl +++ b/OpenCL/m08400_a0.cl @@ -69,8 +69,6 @@ __kernel void m08400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -241,8 +239,6 @@ __kernel void m08400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m08400_a3.cl b/OpenCL/m08400_a3.cl index 769795b39..541cd096f 100644 --- a/OpenCL/m08400_a3.cl +++ b/OpenCL/m08400_a3.cl @@ -67,8 +67,6 @@ __kernel void m08400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -249,8 +247,6 @@ __kernel void m08400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m08900.cl b/OpenCL/m08900.cl index 9d65e1a71..4d3dd3913 100644 --- a/OpenCL/m08900.cl +++ b/OpenCL/m08900.cl @@ -265,8 +265,6 @@ __kernel void m08900_init (__global pw_t *pws, __global const kernel_rule_t *rul const uint4 tmp0 = (uint4) (digest[0], digest[1], digest[2], digest[3]); const uint4 tmp1 = (uint4) (digest[4], digest[5], digest[6], digest[7]); - barrier (CLK_GLOBAL_MEM_FENCE); - tmps[gid].P[k + 0] = tmp0; tmps[gid].P[k + 1] = tmp1; } @@ -331,8 +329,6 @@ __kernel void m08900_comp (__global pw_t *pws, __global const kernel_rule_t *rul for (u32 l = 0; l < SCRYPT_CNT4; l += 4) { - barrier (CLK_GLOBAL_MEM_FENCE); - uint4 tmp; tmp = tmps[gid].P[l + 0]; diff --git a/OpenCL/m09900_a0.cl b/OpenCL/m09900_a0.cl index ff06f126e..7e7459d11 100644 --- a/OpenCL/m09900_a0.cl +++ b/OpenCL/m09900_a0.cl @@ -39,8 +39,6 @@ __kernel void m09900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -104,8 +102,6 @@ __kernel void m09900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m09900_a3.cl b/OpenCL/m09900_a3.cl index 43c464e78..28b1ebe9f 100644 --- a/OpenCL/m09900_a3.cl +++ b/OpenCL/m09900_a3.cl @@ -37,8 +37,6 @@ __kernel void m09900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -108,8 +106,6 @@ __kernel void m09900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m10700.cl b/OpenCL/m10700.cl index a341d3d6e..693691652 100644 --- a/OpenCL/m10700.cl +++ b/OpenCL/m10700.cl @@ -1202,8 +1202,6 @@ __kernel void m10700_loop (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m10800_a0.cl b/OpenCL/m10800_a0.cl index 5e0780118..a84376f81 100644 --- a/OpenCL/m10800_a0.cl +++ b/OpenCL/m10800_a0.cl @@ -39,8 +39,6 @@ __kernel void m10800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -104,8 +102,6 @@ __kernel void m10800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = swap32_S (pws[gid].i[idx]); - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m10800_a3.cl b/OpenCL/m10800_a3.cl index 72ba67211..adff74a7f 100644 --- a/OpenCL/m10800_a3.cl +++ b/OpenCL/m10800_a3.cl @@ -37,8 +37,6 @@ __kernel void m10800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -108,8 +106,6 @@ __kernel void m10800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m11000_a0.cl b/OpenCL/m11000_a0.cl index 6449f5692..cc7310547 100644 --- a/OpenCL/m11000_a0.cl +++ b/OpenCL/m11000_a0.cl @@ -39,8 +39,6 @@ __kernel void m11000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -108,8 +106,6 @@ __kernel void m11000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m11000_a3.cl b/OpenCL/m11000_a3.cl index 5c4788266..b6e7b5f32 100644 --- a/OpenCL/m11000_a3.cl +++ b/OpenCL/m11000_a3.cl @@ -37,8 +37,6 @@ __kernel void m11000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -114,8 +112,6 @@ __kernel void m11000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m11100_a0.cl b/OpenCL/m11100_a0.cl index 2b73633d9..e2b372639 100644 --- a/OpenCL/m11100_a0.cl +++ b/OpenCL/m11100_a0.cl @@ -91,8 +91,6 @@ __kernel void m11100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -258,8 +256,6 @@ __kernel void m11100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m11100_a3.cl b/OpenCL/m11100_a3.cl index 492d4b9af..48216deed 100644 --- a/OpenCL/m11100_a3.cl +++ b/OpenCL/m11100_a3.cl @@ -89,8 +89,6 @@ __kernel void m11100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -292,8 +290,6 @@ __kernel void m11100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m11200_a0.cl b/OpenCL/m11200_a0.cl index c416f9b6a..1ba45efe8 100644 --- a/OpenCL/m11200_a0.cl +++ b/OpenCL/m11200_a0.cl @@ -39,8 +39,6 @@ __kernel void m11200_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -176,8 +174,6 @@ __kernel void m11200_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m11200_a3.cl b/OpenCL/m11200_a3.cl index d162f3690..a76f4d508 100644 --- a/OpenCL/m11200_a3.cl +++ b/OpenCL/m11200_a3.cl @@ -49,8 +49,6 @@ __kernel void m11200_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; @@ -194,8 +192,6 @@ __kernel void m11200_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } sha1_ctx_t ctx0; diff --git a/OpenCL/m11400_a0.cl b/OpenCL/m11400_a0.cl index f75b2e104..72541c78c 100644 --- a/OpenCL/m11400_a0.cl +++ b/OpenCL/m11400_a0.cl @@ -69,8 +69,6 @@ __kernel void m11400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -188,8 +186,6 @@ __kernel void m11400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m11400_a3.cl b/OpenCL/m11400_a3.cl index ec568abd5..cace5a474 100644 --- a/OpenCL/m11400_a3.cl +++ b/OpenCL/m11400_a3.cl @@ -67,8 +67,6 @@ __kernel void m11400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 esalt_len = esalt_bufs[digests_offset].esalt_len; @@ -80,8 +78,6 @@ __kernel void m11400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < esalt_lenv; idx++) { esalt_buf[idx] = esalt_bufs[digests_offset].esalt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; @@ -207,8 +203,6 @@ __kernel void m11400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } const u32 esalt_len = esalt_bufs[digests_offset].esalt_len; @@ -220,8 +214,6 @@ __kernel void m11400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < esalt_lenv; idx++) { esalt_buf[idx] = esalt_bufs[digests_offset].esalt_buf[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } md5_ctx_t ctx0; diff --git a/OpenCL/m11600.cl b/OpenCL/m11600.cl index 44676b478..34b00d893 100644 --- a/OpenCL/m11600.cl +++ b/OpenCL/m11600.cl @@ -177,8 +177,6 @@ __kernel void m11600_loop (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m12400.cl b/OpenCL/m12400.cl index 7ac2b0f0f..1f56c3b4c 100644 --- a/OpenCL/m12400.cl +++ b/OpenCL/m12400.cl @@ -544,8 +544,6 @@ __kernel void m12400_init (__global pw_t *pws, __global const kernel_rule_t *rul for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } u32 tt; diff --git a/OpenCL/m12600_a0.cl b/OpenCL/m12600_a0.cl index 79da6bd9d..22ad327fb 100644 --- a/OpenCL/m12600_a0.cl +++ b/OpenCL/m12600_a0.cl @@ -85,8 +85,6 @@ __kernel void m12600_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -243,8 +241,6 @@ __kernel void m12600_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m12600_a3.cl b/OpenCL/m12600_a3.cl index be4ba4eb2..75dde746a 100644 --- a/OpenCL/m12600_a3.cl +++ b/OpenCL/m12600_a3.cl @@ -83,8 +83,6 @@ __kernel void m12600_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -246,8 +244,6 @@ __kernel void m12600_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m13100_a0.cl b/OpenCL/m13100_a0.cl index d01619666..a3672cff3 100644 --- a/OpenCL/m13100_a0.cl +++ b/OpenCL/m13100_a0.cl @@ -401,8 +401,6 @@ __kernel void m13100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } __local RC4_KEY rc4_keys[64]; @@ -470,8 +468,6 @@ __kernel void m13100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } __local RC4_KEY rc4_keys[64]; diff --git a/OpenCL/m13100_a3.cl b/OpenCL/m13100_a3.cl index 7b1ac5a05..3b7b8b742 100644 --- a/OpenCL/m13100_a3.cl +++ b/OpenCL/m13100_a3.cl @@ -399,8 +399,6 @@ __kernel void m13100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } __local RC4_KEY rc4_keys[64]; @@ -474,8 +472,6 @@ __kernel void m13100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } __local RC4_KEY rc4_keys[64]; diff --git a/OpenCL/m13300_a0.cl b/OpenCL/m13300_a0.cl index 0c3a005b0..a8e85f335 100644 --- a/OpenCL/m13300_a0.cl +++ b/OpenCL/m13300_a0.cl @@ -39,8 +39,6 @@ __kernel void m13300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -106,8 +104,6 @@ __kernel void m13300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m13300_a3.cl b/OpenCL/m13300_a3.cl index 8b69186a2..2ba67e412 100644 --- a/OpenCL/m13300_a3.cl +++ b/OpenCL/m13300_a3.cl @@ -37,8 +37,6 @@ __kernel void m13300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** @@ -110,8 +108,6 @@ __kernel void m13300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { w[idx] = pws[gid].i[idx]; - - barrier (CLK_GLOBAL_MEM_FENCE); } /** diff --git a/OpenCL/m15700.cl b/OpenCL/m15700.cl index a1971093f..8f5de92a9 100644 --- a/OpenCL/m15700.cl +++ b/OpenCL/m15700.cl @@ -396,8 +396,6 @@ __kernel void m15700_init (__global pw_t *pws, __global const kernel_rule_t *rul const uint4 tmp0 = (uint4) (digest[0], digest[1], digest[2], digest[3]); const uint4 tmp1 = (uint4) (digest[4], digest[5], digest[6], digest[7]); - barrier (CLK_GLOBAL_MEM_FENCE); - tmps[gid].P[k + 0] = tmp0; tmps[gid].P[k + 1] = tmp1; } @@ -462,8 +460,6 @@ __kernel void m15700_comp (__global pw_t *pws, __global const kernel_rule_t *rul for (u32 l = 0; l < SCRYPT_CNT4; l += 4) { - barrier (CLK_GLOBAL_MEM_FENCE); - uint4 tmp; tmp = tmps[gid].P[l + 0];