[Autobackout][FuncReg]Revert of change: 7c428db

 Enable Device scope for OpControlBarrier

Invoke global_barrier for Device scope. Remove call to OpControlBarrier in global_barrier to avoid recursion.

Author: ficol

IGC/BiFModule/Implementation/barrier.cl

@@ -124,11 +124,7 @@ void SPIRV_OVERLOADABLE SPIRV_BUILTIN(ControlBarrier, _i32_i32_i32, )(int Execut
         __intel_atomic_work_item_fence( Memory, Semantics );
     }
 
-    if (Execution == Device)
-    {
-        global_barrier();
-    }
-    else  if( Execution <= Workgroup )
+    if( Execution <= Workgroup )
     {
         __builtin_IB_thread_group_barrier();
     }
@@ -286,18 +282,15 @@ void __builtin_spirv_OpMemoryNamedBarrierWrapperOCL_p3__namedBarrier_i32_i32(loc
 }
 
 __global volatile uchar* __builtin_IB_get_sync_buffer();
-uint __intel_get_local_linear_id( void );
-uint __intel_get_local_size( void );
 
 void global_barrier()
 {
     //Make sure each WKG item hit the barrier.
-    __intel_atomic_work_item_fence(Device, AcquireRelease | CrossWorkgroupMemory);
-    __builtin_IB_thread_group_barrier();
+    barrier(CLK_GLOBAL_MEM_FENCE);
 
     __global volatile uchar* syncBuffer = __builtin_IB_get_sync_buffer();
-    bool firstThreadPerWg = __intel_is_first_work_group_item();
-    uint groupLinearId = (__builtin_IB_get_group_id(2) * __builtin_IB_get_num_groups(1) * __builtin_IB_get_num_groups(0)) + (__builtin_IB_get_group_id(1) * __builtin_IB_get_num_groups(0)) + __builtin_IB_get_group_id(0);
+    bool firstThreadPerWg = (get_local_id(0) == 0) && (get_local_id(1) == 0) && (get_local_id(2) == 0);
+    size_t groupLinearId = (get_group_id(2) * get_num_groups(1) * get_num_groups(0)) + (get_group_id(1) * get_num_groups(0)) + get_group_id(0);
 
     //Now first thread of each wkg writes to designated place in syncBuffer
     if (firstThreadPerWg)
@@ -306,27 +299,26 @@ void global_barrier()
         atomic_work_item_fence(CLK_GLOBAL_MEM_FENCE, memory_order_release, memory_scope_device); // == write_mem_fence(CLK_GLOBAL_MEM_FENCE);
     }
 
-    uint numGroups = __builtin_IB_get_num_groups(0) * __builtin_IB_get_num_groups(1) * __builtin_IB_get_num_groups(2);
+    size_t numGroups = get_num_groups(0) * get_num_groups(1) * get_num_groups(2);
     //Higher wkg ids tend to not have work to do in all cases, therefore I choose last wkg to wait for the others, as it is most likely it will hit this code sooner.
     if (groupLinearId == (numGroups - 1))
     {
-        uint localSize = __intel_get_local_size();
+        size_t localSize = get_local_size(0) * get_local_size(1) * get_local_size(2);
         //24 -48 case
         volatile uchar Value;
         do
         {
             atomic_work_item_fence(CLK_GLOBAL_MEM_FENCE, memory_order_acquire, memory_scope_device); // == read_mem_fence(CLK_GLOBAL_MEM_FENCE);
             Value = 1;
-            for (uint i = __intel_get_local_linear_id(); i < numGroups; i += localSize)
+            for (size_t i = get_local_linear_id(); i < numGroups; i += localSize)
             {
                 Value = Value & syncBuffer[i];
             }
 
         } while (Value == 0);
-        __intel_atomic_work_item_fence(Device, AcquireRelease | CrossWorkgroupMemory);
-        __builtin_IB_thread_group_barrier();
+        barrier(CLK_GLOBAL_MEM_FENCE);
 
-        for (uint i = __intel_get_local_linear_id(); i < numGroups; i += localSize)
+        for (size_t i = get_local_linear_id(); i < numGroups; i += localSize)
         {
             syncBuffer[i] = 0;
         }
@@ -339,8 +331,7 @@ void global_barrier()
            atomic_work_item_fence(CLK_GLOBAL_MEM_FENCE, memory_order_acquire, memory_scope_device); // == read_mem_fence(CLK_GLOBAL_MEM_FENCE);
         };
     }
-    __intel_atomic_work_item_fence(Device, AcquireRelease | CrossWorkgroupMemory);
-    __builtin_IB_thread_group_barrier();
+    barrier(CLK_GLOBAL_MEM_FENCE);
 }
 
 void system_memfence(char fence_typed_memory)