Add/restore missing synchronization

Change-Id: Ic1296c58fe984135697430aae66fdcbca84138d0
ROCm · Jan 27, 2020 · c214a58 · c214a58
1 parent 8f441a8
commit c214a58
Showing 1 changed file with 8 additions and 3 deletions.
diff --git a/ockl/src/cg.cl b/ockl/src/cg.cl
@@ -126,16 +126,21 @@ __attribute__((convergent)) void
 __ockl_multi_grid_sync(void)
 {
     __llvm_fence_sc_sys();
+    uint nwm1 = (uint)__ockl_get_num_groups(0) * (uint)__ockl_get_num_groups(1) * (uint)__ockl_get_num_groups(2) - 1;
+    bool cwwi = choose_one_workgroup_workitem();
+
+    if (cwwi)
+        __ockl_gws_barrier(nwm1, 0);
+
+    __builtin_amdgcn_s_barrier();
 
     if (choose_one_grid_workitem()) {
         __constant struct mg_info *m = (__constant struct mg_info *)get_mg_info_arg();
         multi_grid_sync(m->mgs, m->num_grids);
     }
 
-    if (choose_one_workgroup_workitem()) {
-        uint nwm1 = (uint)__ockl_get_num_groups(0) * (uint)__ockl_get_num_groups(1) * (uint)__ockl_get_num_groups(2) - 1;
+    if (cwwi)
         __ockl_gws_barrier(nwm1, 0);
-    }
 
     __builtin_amdgcn_s_barrier();
 }