From 24a2fb01aaf7f761a66ac58ff60809e6566e0ff4 Mon Sep 17 00:00:00 2001 From: jsteube Date: Mon, 7 Aug 2017 18:58:23 +0200 Subject: [PATCH] Fix missing barrier in -m 8500 --- OpenCL/m08500_a0.cl | 2 ++ OpenCL/m08500_a1.cl | 4 ++++ 2 files changed, 6 insertions(+) diff --git a/OpenCL/m08500_a0.cl b/OpenCL/m08500_a0.cl index e81915163..2a1d7a98a 100644 --- a/OpenCL/m08500_a0.cl +++ b/OpenCL/m08500_a0.cl @@ -559,6 +559,8 @@ __kernel void m08500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule s_skb[7][i] = c_skb[7][i]; } + barrier (CLK_LOCAL_MEM_FENCE); + if (gid >= gid_max) return; /** diff --git a/OpenCL/m08500_a1.cl b/OpenCL/m08500_a1.cl index ade03fca5..50d221f52 100644 --- a/OpenCL/m08500_a1.cl +++ b/OpenCL/m08500_a1.cl @@ -557,6 +557,8 @@ __kernel void m08500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule s_skb[7][i] = c_skb[7][i]; } + barrier (CLK_LOCAL_MEM_FENCE); + if (gid >= gid_max) return; /** @@ -711,6 +713,8 @@ __kernel void m08500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule s_skb[7][i] = c_skb[7][i]; } + barrier (CLK_LOCAL_MEM_FENCE); + if (gid >= gid_max) return; /**