xref: /dports/lang/intel-compute-runtime/compute-runtime-21.52.22081/opencl/source/scheduler/scheduler.cl

/*
 * Copyright (C) 2018-2021 Intel Corporation
 *
 * SPDX-License-Identifier: MIT
 *
 */

#ifndef SCHEDULER_EMULATION
#include "device_enqueue.h"
#endif

// float passed as int
extern float __intel__getProfilingTimerResolution();

#ifndef EMULATION_ENTER_FUNCTION
#define EMULATION_ENTER_FUNCTION( )
#endif

#ifndef NULL
#define NULL                                    0
#endif

#define SIMD8                                   0
#define SIMD16                                  1
#define SIMD32                                  2

#define SHORT_SIZE_IN_BYTES                     2
#define DWORD_SIZE_IN_BYTES                     4
#define QWORD_SIZE_IN_BYTES                     8

#define MAX_GLOBAL_ARGS                         255

#define MASK_LOW16_BITS                         0xFFFF
//Currently setting to 8
#define MAX_WALKERS_IN_PARALELL                 PARALLEL_SCHEDULER_HW_GROUPS
//Need 4 uints per walker ( command packet offset, slb offest, dsh offset, idt offset,  + 1 to store total
#define PARALLEL_SCHEDULER_OFFSETS_NUMBER       4
#define PARALLEL_SCHEDULER_LOCAL_MEM_SIZE       ( MAX_WALKERS_IN_PARALELL * PARALLEL_SCHEDULER_OFFSETS_NUMBER + 1 )
//Last index
#define TOTAL_ENQUEUES_FOUND                    ( PARALLEL_SCHEDULER_LOCAL_MEM_SIZE - 1 )

//CURBE STUFF, only entries that really needs to be patched
#define SCHEDULER_DATA_PARAMETER_KERNEL_ARGUMENT                                    1
#define SCHEDULER_DATA_PARAMETER_LOCAL_WORK_SIZE                                    2
#define SCHEDULER_DATA_PARAMETER_GLOBAL_WORK_SIZE                                   3
#define SCHEDULER_DATA_PARAMETER_NUM_WORK_GROUPS                                    4
#define SCHEDULER_DATA_PARAMETER_WORK_DIMENSIONS                                    5
#define SCHEDULER_DATA_PARAMETER_SUM_OF_LOCAL_MEMORY_OBJECT_ARGUMENT_SIZES          8
#define SCHEDULER_DATA_PARAMETER_GLOBAL_WORK_OFFSET                                 16
#define SCHEDULER_DATA_PARAMETER_NUM_HARDWARE_THREADS                               17
#define SCHEDULER_DATA_PARAMETER_PARENT_EVENT                                       22
#define SCHEDULER_DATA_PARAMETER_ENQUEUED_LOCAL_WORK_SIZE                           28

#define SCHEDULER_DATA_PARAMETER_IMAGE_WIDTH                                        ( 9 + SCHEDULER_DATA_PARAMETER_IMAGES_CURBE_SHIFT )
#define SCHEDULER_DATA_PARAMETER_IMAGE_HEIGHT                                       ( 10 + SCHEDULER_DATA_PARAMETER_IMAGES_CURBE_SHIFT )
#define SCHEDULER_DATA_PARAMETER_IMAGE_DEPTH                                        ( 11 + SCHEDULER_DATA_PARAMETER_IMAGES_CURBE_SHIFT )
#define SCHEDULER_DATA_PARAMETER_IMAGE_CHANNEL_DATA_TYPE                            ( 12 + SCHEDULER_DATA_PARAMETER_IMAGES_CURBE_SHIFT )
#define SCHEDULER_DATA_PARAMETER_IMAGE_CHANNEL_ORDER                                ( 13 + SCHEDULER_DATA_PARAMETER_IMAGES_CURBE_SHIFT )
#define SCHEDULER_DATA_PARAMETER_IMAGE_ARRAY_SIZE                                   ( 18 + SCHEDULER_DATA_PARAMETER_IMAGES_CURBE_SHIFT )
#define SCHEDULER_DATA_PARAMETER_IMAGE_NUM_SAMPLES                                  ( 20 + SCHEDULER_DATA_PARAMETER_IMAGES_CURBE_SHIFT )
#define SCHEDULER_DATA_PARAMETER_IMAGE_NUM_MIP_LEVELS                               ( 27 + SCHEDULER_DATA_PARAMETER_IMAGES_CURBE_SHIFT )
#define SCHEDULER_DATA_PARAMETER_IMAGE_OBJECT_ID                                    ( 35 + SCHEDULER_DATA_PARAMETER_IMAGES_CURBE_SHIFT )
#define SCHEDULER_DATA_PARAMETER_IMAGE_SRGB_CHANNEL_ORDER                           ( 39 + SCHEDULER_DATA_PARAMETER_IMAGES_CURBE_SHIFT )

#define DATA_PARAMETER_SAMPLER_ADDRESS_MODE                                         ( 14 + SCHEDULER_DATA_PARAMETER_SAMPLER_ADDED_VALUE )
#define DATA_PARAMETER_SAMPLER_NORMALIZED_COORDS                                    ( 15 + SCHEDULER_DATA_PARAMETER_SAMPLER_ADDED_VALUE )
#define DATA_PARAMETER_SAMPLER_COORDINATE_SNAP_WA_REQUIRED                          ( 21 + SCHEDULER_DATA_PARAMETER_SAMPLER_ADDED_VALUE )
#define SCHEDULER_DATA_PARAMETER_SAMPLER_OBJECT_ID                                  ( 35 + SCHEDULER_DATA_PARAMETER_SAMPLER_ADDED_VALUE )

//CURBE STUFF, only entries that really needs to be patched
#define SCHEDULER_DATA_PARAMETER_KERNEL_ARGUMENT_MASK                               ( 1 << 1 )
#define SCHEDULER_DATA_PARAMETER_LOCAL_WORK_SIZE_MASK                               ( 1 << 2 )
#define SCHEDULER_DATA_PARAMETER_GLOBAL_WORK_SIZE_MASK                              ( 1 << 3 )
#define SCHEDULER_DATA_PARAMETER_NUM_WORK_GROUPS_MASK                               ( 1 << 4 )
#define SCHEDULER_DATA_PARAMETER_WORK_DIMENSIONS_MASK                               ( 1 << 5 )
#define SCHEDULER_DATA_PARAMETER_SUM_OF_LOCAL_MEMORY_OBJECT_ARGUMENT_SIZES_MASK     ( 1 << 8 )
#define SCHEDULER_DATA_PARAMETER_GLOBAL_WORK_OFFSET_MASK                            ( 1 << 16 )
#define SCHEDULER_DATA_PARAMETER_NUM_HARDWARE_THREADS_MASK                          ( 1 << 17 )
#define SCHEDULER_DATA_PARAMETER_PARENT_EVENT_MASK                                  ( 1 << 22 )
#define SCHEDULER_DATA_PARAMETER_ENQUEUED_LOCAL_WORK_SIZE_MASK                      ( 1 << 28 )
#define SCHEDULER_DATA_PARAMETER_IMAGE_CURBE_ENTRIES                                ( ( ulong ) 1 << SCHEDULER_DATA_PARAMETER_IMAGES_CURBE_SHIFT )
#define SCHEDULER_DATA_PARAMETER_GLOBAL_POINTER                                     ( ( ( ulong ) 1 ) << SCHEDULER_DATA_PARAMETER_GLOBAL_POINTER_SHIFT )


#define SCHEDULER_DATA_PARAMETER_SAMPLER_MASK                                       ( ( ( ulong ) 1 ) << SCHEDULER_DATA_PARAMETER_SAMPLER_SHIFT )
//Error codes
#define SCHEDULER_CURBE_TOKEN_MISSED                                                10
#define SCHEDULER_CURBE_ARGUMENTS_SIZE_MISMATCH                                     11

#define CAN_BE_RECLAIMED                                                            123456

#define SCHEDULER_MSF_INITIAL                                                       1
#define SCHEDULER_MSF_SECOND                                                        2

//Uncomment to enable logging debug data
//#define    ENABLE_DEBUG_BUFFER    1

#ifdef ENABLE_DEBUG_BUFFER
//Update DebugDataInfo types in device_enqueue.h and PrintDebugDataBuffer() in cmd_queue_device.cpp

//Flags
#define    DDB_HAS_DATA_INFO        ( 1 << 0 )
#define    DDB_SCHEDULER_PROFILING  ( 1 << 1 )

#define    DDB_ALL                  ( 0xffffffff )

#endif

//Turn this to 1 to turn on debug calls, notice that it will cause up to 10 x longer time to build scheduler
//#define    SCHEDULER_DEBUG_MODE 1

//#define    DISABLE_RESOURCE_RECLAMATION                                                 1
/*
Resource reclamation procedure
1. Move all new command packets from queue_t to qstorage
2. In case there is place in storage for whole queue, reclaim space on queue
3. Construct stack basing on new commands added in the qstorage
4. Browse stack to find next item for execution
5. When you take item from the stack and schedule it , reclaim place on qstorage buffer
*/

typedef struct
{
    uint3           ActualLocalSize;
    uint3           WalkerDimSize;
    uint3           WalkerStartPoint;
} IGIL_WalkerData;

typedef struct
{
    uint3           LocalWorkSize;
    uint3           TotalDimSize;
    IGIL_WalkerData WalkerArray[ 8 ];
} IGIL_WalkerEnumeration;

inline void patchDword( __global uint* patchedDword, uint startOffset, uint endOffset, uint value )
{
    uint LeftMask  = ALL_BITS_SET_DWORD_MASK >> ( DWORD_SIZE_IN_BITS - endOffset - 1 );
    uint RightMask = ALL_BITS_SET_DWORD_MASK << ( startOffset );
    uint CleanMask = ~( RightMask & LeftMask );
    *patchedDword &= CleanMask;
    *patchedDword |= ( value << startOffset );
}

inline __global IGIL_KernelAddressData* IGIL_GetKernelAddressData( __global IGIL_KernelDataHeader* pKernelReflection, uint blockId )
{
    return ( __global IGIL_KernelAddressData* ) ( &pKernelReflection->m_data[ blockId ] );
}
__global IGIL_KernelData* IGIL_GetKernelData( __global IGIL_KernelDataHeader* pKernelReflection, uint blockId )
{
    __global IGIL_KernelAddressData* pKernelAddressData = IGIL_GetKernelAddressData( pKernelReflection, blockId );
    uint Offset                                         = pKernelAddressData->m_KernelDataOffset;
    __global char* pKernelReflectionRaw                 = ( __global char * ) pKernelReflection;
    return ( __global IGIL_KernelData* ) ( &pKernelReflectionRaw[ Offset ] );
}


inline __global IGIL_CommandHeader* TEMP_IGIL_GetCommandHeader( __global IGIL_CommandQueue* q, uint offset )
{
    __global uchar *pQueueRaw = (__global uchar *) q;

    __global IGIL_CommandHeader* pCommand = ( __global IGIL_CommandHeader* )( pQueueRaw + offset );

    return pCommand;
}
//Make sure enough command packets are in command queue before calling this function.
__global IGIL_CommandHeader* TEMP_IGIL_GetNthCommandHeader( __global IGIL_CommandQueue* q, uint initialOffset, uint number )
{
    __global uchar *pQueueRaw = (__global uchar *) q;

    __global IGIL_CommandHeader* pCommand = ( __global IGIL_CommandHeader* )( pQueueRaw + initialOffset );
    uint Offset = initialOffset;
    //Traverse queue_t unless nth command packet is found
    while( number > 0 )
    {
        Offset  += pCommand->m_commandSize;
        pCommand = TEMP_IGIL_GetCommandHeader( q, Offset );
        number--;
    }

    return pCommand;
}

//Make sure enough command packets are in command queue before calling this function.
uint TEMP_IGIL_GetNthCommandHeaderOffset( __global IGIL_CommandQueue* q, uint initialOffset, uint number )
{
    __global uchar *pQueueRaw = (__global uchar *) q;

    __global IGIL_CommandHeader* pCommand = ( __global IGIL_CommandHeader* )( pQueueRaw + initialOffset );
    uint Offset = initialOffset;
    //Traverse queue_t unless nth command packet is found
    while( number > 0 )
    {
        Offset  += pCommand->m_commandSize;
        pCommand = TEMP_IGIL_GetCommandHeader( q, Offset );
        number--;
    }

    return Offset;
}

__global IGIL_CommandHeader* GetCommandHeaderFromStorage( __global uint* queueStorage, uint offset )
{
    __global uchar *pQueueRaw = ( __global uchar * ) queueStorage;

    __global IGIL_CommandHeader* pCommand = ( __global IGIL_CommandHeader* )( pQueueRaw + offset );

    return pCommand;
}

inline queue_t TEMP_IGIL_GetQueueT( IGIL_CommandHeader * queuePtr )
{
    return __builtin_astype( queuePtr, queue_t );
}

inline __global IGIL_DeviceEvent* TEMP_IGIL_GetDeviceEvents( __global IGIL_EventPool *pool )
{
    return ( __global IGIL_DeviceEvent * )( pool + 1 );
}

inline __global IGIL_DeviceEvent* TEMP_IGIL_GetDeviceEvent( __global IGIL_EventPool *pool, uint eventId )
{
    __global IGIL_DeviceEvent * pEvent = ( __global IGIL_DeviceEvent * )( pool + 1 );

    return ( __global IGIL_DeviceEvent * )( pEvent + eventId );
}
bool SetEventState( __global IGIL_EventPool *pool, uint eventId, int state )
{
    __global IGIL_DeviceEvent* pDeviceEvent = TEMP_IGIL_GetDeviceEvent( pool, eventId );
    pDeviceEvent->m_state = state;
    return true;
}

void TEMP_IGIL_FreeEvent( clk_event_t event, __global IGIL_EventPool *pool )
{
    //Offset into the event data in the pool
    __global IGIL_DeviceEvent *events = TEMP_IGIL_GetDeviceEvents( pool );

    atomic_xchg( &events[ ( uint )(size_t)__builtin_astype( event, void* ) ].m_state, IGIL_EVENT_UNUSED );
}

void  IGILLOCAL_MEMCPY_GTOG( __global void* pDst, __global void* pSrc, int numBytes )
{
    numBytes = numBytes >> 2;
    for( int i = 0; i < numBytes; i++ )
    {
        ( ( __global uint* )pDst ) [ i ] = ( ( __global int* )pSrc )[ i ];
    }
}

//Global memcpy running on all witems possible, make sure it's run from all hw threads.
void GLOBAL_MEMCPYUINT( __global void* pDst, __global void* pSrc, int numBytes )
{
    uint total_local_size   = get_local_size( 0 );
    uint LoopCtr            = numBytes / ( total_local_size * DWORD_SIZE_IN_BYTES );
    uint LeftOver           = numBytes % ( total_local_size * DWORD_SIZE_IN_BYTES );
    uint Lid                = get_local_id( 0 );
    uint i                  = 0;

    //Main copy
    for( i = 0; i < LoopCtr; i++ )
    {
         ( ( __global uint* )pDst ) [ Lid + total_local_size * i ] = ( ( __global uint* )pSrc )[ Lid + total_local_size * i ];
    }
    //Copy what's left
    if( LeftOver != 0 )
    {
        if( Lid * DWORD_SIZE_IN_BYTES < LeftOver )
        {
            ( ( __global uint* )pDst ) [ Lid + total_local_size * i ] = ( ( __global uint* )pSrc )[ Lid + total_local_size * i ];
        }
    }
}

//In SIMD8 to fully use cachelines copy in portions of uint2 , 8 bytes x 8 witems = 64 bytes cacheline size.
//Global memcpy running on all witems possible, make sure it's run from all hw threads.
void GLOBAL_MEMCPY( __global void* pDst, __global void* pSrc, int numBytes )
{
    //In case I need dword copy use uint version of this function.
    if( ( numBytes % ( DWORD_SIZE_IN_BYTES * 2 ) ) != 0 )
    {
       GLOBAL_MEMCPYUINT( pDst, pSrc, numBytes );
    }
    else
    {
        uint total_local_size   = get_local_size( 0 );
        uint LoopCtr            = numBytes / ( total_local_size * DWORD_SIZE_IN_BYTES * 2 );
        uint LeftOver           = numBytes % ( total_local_size * DWORD_SIZE_IN_BYTES * 2 );
        uint Lid                = get_local_id( 0 );
        uint i                  = 0;

        //Main copy
        for( i = 0; i < LoopCtr; i++ )
        {
             ( ( __global uint2* )pDst ) [ Lid + total_local_size * i ] = ( ( __global uint2* )pSrc )[ Lid + total_local_size * i ];
        }
        //Copy what's left
        if( LeftOver != 0 )
        {
            if( Lid * DWORD_SIZE_IN_BYTES * 2 < LeftOver )
            {
                ( ( __global uint2* )pDst ) [ Lid + total_local_size * i ] = ( ( __global uint2* )pSrc )[ Lid + total_local_size * i ];
            }
        }
    }
}

//This works only for 32 bit types
uint GetNextPowerof2( uint number )
{
    number --;
    number |= number >> 1;
    number |= number >> 2;
    number |= number >> 4;
    number |= number >> 8;
    number |= number >> 16;
    number ++;
    return number;
}

#define OCLRT_ALIGN( a, b )                 ( ( ( ( a ) % ( b ) ) != 0 ) ?  ( ( a ) - ( ( a ) % ( b ) ) + ( b ) ) : ( a ) )
#define OCLRT_MAX( a, b )                   ( ( ( a ) > ( b ) ) ?  ( a ) : ( b ) )

#ifndef SCHEDULER_EMULATION
#include "scheduler_definitions.h"
#endif

#ifdef ENABLE_DEBUG_BUFFER
//Adds private uint* data to debug buffer
//  ddb - global buffer to keep data
//  src - source data
//  numberOfElements
//  localID - WI ID , when 0xffffffff all WI copy data
//  returns 0 when data was copied, else -1
int AddToDebugBufferParallel( __global DebugDataBuffer* ddb, uint* pSrc, uint numberOfElements, uint localID )
{
    if( ddb->m_flags == 0 )
    {
        //All work items in local group copies data
        if( localID == 0xffffffff )
        {
            //Check if there is enough place for new data in ddb for every workitem in local group
            if( ( ( ddb->m_size / sizeof( uint ) - ddb->m_offset ) * get_local_size(0) ) >= numberOfElements )
            {
                uint startIndex = atomic_add( &ddb->m_offset, numberOfElements );
                uint i;
                int srcPos;
                for( i = startIndex, srcPos = 0; i < startIndex + numberOfElements; i++, srcPos++ )
                {
                    ddb->m_data[ i ] = pSrc[ srcPos ];
                }
                return 0;
            }
        }
        else
        {
            //Check if there is enough place for new data in ddb for one workitem in local group
            if( ( ddb->m_size / sizeof( uint ) - ddb->m_offset ) >= numberOfElements )
            {
                if( get_local_id(0) == localID )
                {
                    uint startIndex = atomic_add( &ddb->m_offset, numberOfElements );
                    uint i;
                    int srcPos;
                    for( i = startIndex, srcPos = 0; i < startIndex + numberOfElements; i++, srcPos++ )
                    {
                        ddb->m_data[ i ] = pSrc[ srcPos ];
                    }
                    return 0;
                }
            }
        }
    }
    return -1;
}

//Adds private uint to debug buffer
int AddToDBParallel(__global DebugDataBuffer* ddb, uint pSrc, uint localID)
{
    return AddToDebugBufferParallel(ddb, &pSrc, 1, localID);
}

//Adds global uint* data to debug buffer
//  ddb - global buffer to keep data
//  src - source data
//  numberOfElements
//  localID - WI ID , when 0xffffffff all WI copy data
//  returns 0 when data was copied, else -1
int AddGlobalToDebugBufferParallel( __global DebugDataBuffer* ddb, __global uint* pSrc, uint numberOfElements, uint localID )
{
    if( ddb->m_flags == 0 )
    {
        //All work items in local group copies data
        if( localID == 0xffffffff )
        {
            //Check if there is enough place for new data in ddb for every workitem in local group
            if( ( ( ddb->m_size / sizeof( uint ) - ddb->m_offset ) * get_local_size(0) ) >= numberOfElements )
            {
                uint startIndex = atomic_add( &ddb->m_offset, numberOfElements );
                uint i;
                int srcPos;
                for( i = startIndex, srcPos = 0; i < startIndex + numberOfElements; i++, srcPos++ )
                {
                    ddb->m_data[ i ] = pSrc[ srcPos ];
                }
                return 0;
            }
        }
        else
        {
            //Check if there is enough place for new data in ddb for one workitem in local group
            if( ( ddb->m_size / sizeof( uint ) - ddb->m_offset ) >= numberOfElements )
            {
                if( get_local_id(0) == localID )
                {
                    uint startIndex = atomic_add( &ddb->m_offset, numberOfElements );
                    uint i;
                    int srcPos;
                    for( i = startIndex, srcPos = 0; i < startIndex + numberOfElements; i++, srcPos++ )
                    {
                        ddb->m_data[ i ] = pSrc[ srcPos ];
                    }
                    return 0;
                }
            }
        }
    }
    return -1;
}



//Adds private uint data to debug buffer
//  ddb - global buffer to keep data
//  src - source data
//  dataType - enum defining added data type
//  returns 0 when data was copied, else -1
int AddToDebugBuffer( __global DebugDataBuffer* ddb, __private ulong src, uint dataType, uint localID )
{
    if( ddb->m_flags == 0 || ddb->m_flags == DDB_HAS_DATA_INFO )
    {
        //All work items in local group copies data
        if( localID == 0xffffffff )
        {
            //Check if there is enough place for new data in ddb
            if( ( ( ddb->m_stackTop - ddb->m_dataInfoTop ) >= 4 * get_local_size( 0 ) ) && ( ddb->m_dataInfoTop > ddb->m_stackTop ) )
            {
                //Check flags
                if( ddb->m_flags == 0 || ddb->m_flags == DDB_HAS_DATA_INFO )
                {
                    uint startIndex                 = atomic_add( &ddb->m_offset, 1 );
                    uint dataIndex                  = atomic_sub( &ddb->m_dataInfoTop, ddb->m_dataInfoSize );
                    uint stackTop                   = atomic_add( &ddb->m_stackTop, 8 );

                    __global uchar* pCharDebugQueue = ( __global uchar * )ddb;
                    __global uint* dest             = ( __global uint* )&pCharDebugQueue[ stackTop ];
                    __global DebugDataInfo* debugInfo;

                    dest[ 0 ]                 = ( uint )src & 0xffffffff;
                    dest[ 1 ]               = ( src & 0xffffffff00000000 ) >> 32;
                    debugInfo               = ( __global DebugDataInfo* )( &pCharDebugQueue[ dataIndex ] );
                    debugInfo->m_dataType   = ( DebugDataTypes )dataType;
                    debugInfo->m_dataSize   = 8;
                    ddb->m_flags            |= DDB_HAS_DATA_INFO;
                    return 0;
                }
            }
        }
        else
        {
            //Check if there is enough place for new data in ddb
            if( ( ( ddb->m_stackTop - ddb->m_dataInfoTop ) >= 8 ) && ( ddb->m_dataInfoTop > ddb->m_stackTop ) )
            {
                if( get_local_id( 0 ) == localID )
                {
                    uint startIndex                 = atomic_add( &ddb->m_offset, 1 );
                    uint dataIndex                  = atomic_sub( &ddb->m_dataInfoTop, ddb->m_dataInfoSize );
                    uint stackTop                   = atomic_add( &ddb->m_stackTop, 8 );

                    __global uchar* pCharDebugQueue = ( __global uchar * )ddb;
                    __global uint* dest             = ( __global uint* )&pCharDebugQueue[ stackTop ];
                    __global DebugDataInfo* debugInfo;

                    dest[ 0 ]                 = ( uint )src & 0xffffffff;
                    dest[ 1 ]               = ( src & 0xffffffff00000000 ) >> 32;
                    debugInfo               = ( __global DebugDataInfo* )( &pCharDebugQueue[ dataIndex ] );
                    debugInfo->m_dataType   = ( DebugDataTypes )dataType;
                    debugInfo->m_dataSize   = 8;
                    ddb->m_flags            |= DDB_HAS_DATA_INFO;
                    return 0;
                }

            }
        }
    }
    return -1;
}

//Adds data to debug buffer
//  ddb - global buffer to keep data
//  src - source
//  bytes - number of bytes from src to put into ddb
//  dataType - enum defining added data type
//  returns 0 when data was copied, else -1
int AddGlobalToDebugBuffer( __global DebugDataBuffer* ddb, __global uchar* src, uint bytes, uint dataType )
{
    if( get_local_id( 0 ) )
    {
        //Check if there is enough place for new data in ddb
        if( ( ( ddb->m_stackTop - ddb->m_dataInfoTop ) >= bytes ) && ( ddb->m_dataInfoTop > ddb->m_stackTop ) )
        {
            //Check flags
            if( ddb->m_flags == 0 || ddb->m_flags == DDB_HAS_DATA_INFO )
            {
                __global uchar* pCharDebugQueue = ( __global uchar * )ddb;
                __global DebugDataInfo* debugInfo;
                IGILLOCAL_MEMCPY_GTOG( ( &pCharDebugQueue[ ddb->m_stackTop ] ), ( src ), ( int )bytes );
                debugInfo = ( __global DebugDataInfo* )( &pCharDebugQueue[ ddb->m_dataInfoTop ] );
                debugInfo->m_dataType = ( DebugDataTypes )dataType;
                debugInfo->m_dataSize = bytes;

                ddb->m_dataInfoTop  = ddb->m_dataInfoTop - ddb->m_dataInfoSize;
                ddb->m_stackTop     = ddb->m_stackTop + bytes;
                ddb->m_offset       = ddb->m_offset + ( bytes / 4 );
                ddb->m_flags        |= DDB_HAS_DATA_INFO;
                return 0;
            }
        }
    }
    return -1;
}

int AddGlobalToDebugBufferAllIds( __global DebugDataBuffer* ddb, __global uchar* src, uint bytes, uint dataType, uint localID )
{
    if( ddb->m_flags == 0 || ddb->m_flags == DDB_HAS_DATA_INFO )
    {
        //Check if there is enough place for new data in ddb
        if( ( ( ddb->m_stackTop - ddb->m_dataInfoTop ) >= bytes ) && ( ddb->m_dataInfoTop > ddb->m_stackTop ) )
        {
            //Check flags
            if( ddb->m_flags == 0 || ddb->m_flags == DDB_HAS_DATA_INFO )
            {
                __global uchar* pCharDebugQueue = ( __global uchar * )ddb;
                __global DebugDataInfo* debugInfo;
                IGILLOCAL_MEMCPY_GTOG( ( &pCharDebugQueue[ ddb->m_stackTop ] ), ( src ), ( int )bytes );
                debugInfo = ( __global DebugDataInfo* )( &pCharDebugQueue[ ddb->m_dataInfoTop ] );
                debugInfo->m_dataType = ( DebugDataTypes )dataType;
                debugInfo->m_dataSize = bytes;

                ddb->m_dataInfoTop  = ddb->m_dataInfoTop - ddb->m_dataInfoSize;
                ddb->m_stackTop     = ddb->m_stackTop + bytes;
                ddb->m_offset       = ddb->m_offset + ( bytes / 4 );
                ddb->m_flags        |= DDB_HAS_DATA_INFO;
                return 0;
            }
        }
    }
    return -1;
}

#endif

#define MAX_SLB_OFFSET ( SECOND_LEVEL_BUFFER_SPACE_FOR_EACH_ENQUEUE * SECOND_LEVEL_BUFFER_NUMBER_OF_ENQUEUES )

#ifndef SCHEDULER_EMULATION
#include "scheduler_builtin_kernel.inl"
#endif

//SOME COMMON CODE FUNCTIONS
//COMMON CODE STARTS HERE
//Not thread safe - make sure it's called in thread safe fashion.

void patchIDData( __global char* dsh,
                  uint blockId,
                  uint numberOfHwThreads,
                  uint slmSize )
{
    __global char* DSHIData                      = ( __global char* )( dsh + SIZEOF_COLOR_CALCULATOR_STATE + ( ( blockId + 1 ) * SIZEOF_INTERFACE_DESCRIPTOR_DATA ) );
    __global uint* DSHIntData                    = ( __global uint* )( DSHIData );
    //Barrier enable is pre-patched on the host.
    patchDword( ( &DSHIntData[ INTERFACE_DESCRIPTOR_HWTHREADS_NUMBER_DWORD ] ), 0, INTERFACE_DESCRIPTOR_HWTHREADS_UPPER_BIT, numberOfHwThreads );

    //Patch SLM.
    uint SLMPatchValue = GetPatchValueForSLMSize( slmSize );
    patchDword( ( &DSHIntData[ INTERFACE_DESCRIPTOR_HWTHREADS_NUMBER_DWORD ] ), 16, 20, SLMPatchValue );
}
/*
this is how it works :
When constructing primary batch, first IDT table is also constructed, for all blocks, it is constructed as follows:
[0] - parent id
[1 .. x ] block id
[last aligned ] scheduler

now when we enter SLB, we forgot about first IDT, and we point all interface descriptor loads to point at scheduler which was last in the first IDT, to be first in the new IDT.

This way we can copy Interface Descriptors for blocks from the first IDT and assign Interface Descriptors dynamically in scheduler.
*/

void CopyAndPatchIDData( __global char* dsh,
                         uint blockId,
                         uint numberOfHwThreads,
                         uint slmSize,
                         uint interfaceDescriptorOffset,
                         uint blockStartId )
{
    __global char* DSHIData                      = ( __global char* )( dsh + SIZEOF_COLOR_CALCULATOR_STATE + ( ( blockId + blockStartId ) * SIZEOF_INTERFACE_DESCRIPTOR_DATA ) );
    __global uint* DSHIntData                    = ( __global uint* )( DSHIData );

    //Copy to ID InterfaceDescriptorOffset
    __global char* DSHDestIData                  = ( __global char* )( dsh + SIZEOF_COLOR_CALCULATOR_STATE + ( ( IDT_BREAKDOWN + interfaceDescriptorOffset ) * SIZEOF_INTERFACE_DESCRIPTOR_DATA ) );
    __global uint* DSHDestIntData                = ( __global uint* )( DSHDestIData );
    __global uint* DSHDestIntStartData           = DSHDestIntData;

    for( int i = 0; i < ( SIZEOF_INTERFACE_DESCRIPTOR_DATA / 4 ); i++ )
    {
        DSHDestIntData[ i ] = DSHIntData[ i ];
    }

    //Barrier enable is pre-patched on the host.
    patchDword( ( &DSHDestIntStartData[ INTERFACE_DESCRIPTOR_HWTHREADS_NUMBER_DWORD ] ), 0, INTERFACE_DESCRIPTOR_HWTHREADS_UPPER_BIT, numberOfHwThreads );

    //Patch SLM.
    uint SLMPatchValue = GetPatchValueForSLMSize( slmSize );
    patchDword( ( &DSHDestIntStartData[ INTERFACE_DESCRIPTOR_HWTHREADS_NUMBER_DWORD ] ), 16, 20, SLMPatchValue );
}

void CopyAndPatchIDData20( __global char* dsh,
                           uint blockId,
                           uint numberOfHwThreads,
                           uint slmSize,
                           uint interfaceDescriptorOffset,
                           uint blockStartId,
                           uint bToffset,
                           uint dshOffset,
                           uint numOfSamplers
#ifdef ENABLE_DEBUG_BUFFER
                           , __global DebugDataBuffer* DebugQueue
#endif
                           )
{
    EMULATION_ENTER_FUNCTION( );

    __global char* DSHIData                      = ( __global char* )( dsh + SIZEOF_COLOR_CALCULATOR_STATE + ( ( blockId + blockStartId ) * SIZEOF_INTERFACE_DESCRIPTOR_DATA ) );
    __global uint* DSHIntData                    = ( __global uint* )( DSHIData );

    //Copy to ID InterfaceDescriptorOffset
    __global char* DSHDestIData                  = ( __global char* )( dsh + SIZEOF_COLOR_CALCULATOR_STATE + ( ( IDT_BREAKDOWN + interfaceDescriptorOffset ) * SIZEOF_INTERFACE_DESCRIPTOR_DATA ) );
    __global uint* DSHDestIntData                = ( __global uint* )( DSHDestIData );
    __global uint* DSHDestIntStartData           = DSHDestIntData;

    for( int i = 0; i < ( SIZEOF_INTERFACE_DESCRIPTOR_DATA / 4 ); i++ )
    {
        DSHDestIntData[ i ] = DSHIntData[ i ];
    }

    //Barrier enable is pre-patched on the host.
    patchDword( ( &DSHDestIntStartData[ INTERFACE_DESCRIPTOR_HWTHREADS_NUMBER_DWORD ] ), 0, INTERFACE_DESCRIPTOR_HWTHREADS_UPPER_BIT, numberOfHwThreads );

    //Patch BT offset
    patchDword( ( &DSHDestIntStartData[ INTERFACE_DESCRIPTOR_BINDING_TABLE_POINTER_DWORD ] ), 5, 15, ( bToffset >> 5 ) );

    //Patch SLM.
    uint PatchValue = GetPatchValueForSLMSize( slmSize );
    patchDword( ( &DSHDestIntStartData[ INTERFACE_DESCRIPTOR_HWTHREADS_NUMBER_DWORD ] ), 16, 20, PatchValue );

    PatchValue = ( DSHDestIntStartData[ INTERFACE_DESCRIPTOR_SAMPLER_STATE_TABLE_DWORD ] & 0xffffffe0 ) + ( dshOffset );
    patchDword( ( &DSHDestIntStartData[ INTERFACE_DESCRIPTOR_SAMPLER_STATE_TABLE_DWORD ] ), 5, 31, ( ( PatchValue ) >> 5 ) );

    //Samplers in multiple of 4
    numOfSamplers = ( numOfSamplers + 3 ) / 4;
    patchDword( ( &DSHDestIntStartData[ INTERFACE_DESCRIPTOR_SAMPLER_STATE_TABLE_DWORD ] ), 2, 4, numOfSamplers );
}


void patchGpGpuWalker(
    uint secondLevelBatchOffset,
    __global uint* secondaryBatchBuffer,
    uint interfaceDescriptorOffset,
    uint simdSize,
    uint totalLocalWorkSize,
    uint3 dimSize,
    uint3 startPoint,
    uint numberOfHwThreadsPerWg,
    uint indirectPayloadSize,
    uint ioHoffset )
{
    EMULATION_ENTER_FUNCTION( );

    //SlbOffset is expressed in bytes and for cmd it is needed to convert it to dwords
    uint CmdPacketStart = secondLevelBatchOffset / DWORD_SIZE_IN_BYTES;
    //INTERFACE_DESCRIPTOR for GPGPU_WALKER
    //INTERFACE DESCRIPTOR is one plus the block id
    uint PatchOffset = CmdPacketStart + GPGPU_WALKER_INTERFACE_DESCRIPTOR_ID_OFFSET;
    //Patch id data
    patchDword( &( secondaryBatchBuffer[ PatchOffset ] ),
                0, 5, ( interfaceDescriptorOffset ) );
    PatchOffset = CmdPacketStart + GPGPU_WALKER_THREAD_WIDTH_DWORD;
    //THREAD_WIDTH for GPGPU_WALKER
    patchDword( &( secondaryBatchBuffer[ PatchOffset ] ),
                0, 5, ( numberOfHwThreadsPerWg - 1 ) );

    PatchOffset = CmdPacketStart + GPGPU_WALKER_SIMDSIZE_DWORD;

    //SIMD SIZE for GPGPU_WALKER
    //Double Check the bits for SIMDSize
    if( simdSize == 8 )
    {
        patchDword( &( secondaryBatchBuffer[ PatchOffset ] ),
                    30, 31, SIMD8 );
    }
    else if ( simdSize == 16 )
    {
        patchDword( &( secondaryBatchBuffer[ PatchOffset ] ),
                    30, 31, SIMD16 );
    }
    else
    {
        patchDword( &( secondaryBatchBuffer[ PatchOffset ] ),
                    30, 31, SIMD32 );
    }

    //XDIM for GPGPU_WALKER
    secondaryBatchBuffer[ CmdPacketStart + GPGPU_WALKER_XDIM_DWORD ] = dimSize.x;
    secondaryBatchBuffer[ CmdPacketStart + GPGPU_WALKER_GROUP_ID_START_X ] = startPoint.x;
    //YDIM
    secondaryBatchBuffer[ CmdPacketStart + GPGPU_WALKER_YDIM_DWORD ] = dimSize.y;
    secondaryBatchBuffer[ CmdPacketStart + GPGPU_WALKER_GROUP_ID_START_Y ] = startPoint.y;
    //ZDIM for GPGPU_WALKER
    secondaryBatchBuffer[ CmdPacketStart + GPGPU_WALKER_ZDIM_DWORD ] = dimSize.z;
    secondaryBatchBuffer[ CmdPacketStart + GPGPU_WALKER_GROUP_ID_START_Z ] = startPoint.z;

    //XMASK for GPGPU_WALKER
    uint mask = ( 1 << ( totalLocalWorkSize % simdSize ) ) - 1;
    if( mask == 0 )
        mask = ~0;

    secondaryBatchBuffer[ CmdPacketStart + GPGPU_WALKER_XMASK_DWORD ] = mask;

    //YMASK for GPGPU_WALKER
    uint YMask = ~0;

    secondaryBatchBuffer[ CmdPacketStart + GPGPU_WALKER_YMASK_DWORD ] = YMask;


    patchDword( &( secondaryBatchBuffer[ CmdPacketStart + GPGPU_WALKER_INDIRECT_DATA_LENGTH_OFFSET ] ),
                0, 16, indirectPayloadSize );

    patchDword( &( secondaryBatchBuffer[ CmdPacketStart + GPGPU_WALKER_INDIRECT_START_ADDRESS_OFFSET ] ),
                0, 31, ioHoffset );
}

int PatchMediaStateFlush(
    uint secondLevelBatchOffset,
    __global uint* secondaryBatchBuffer,
    uint interfaceDescriptorOffset,
    uint msfNumber )
{
    //SlbOffset is expressed in bytes and for cmd it is needed to convert it to dwords
    uint CmdPacketStart = secondLevelBatchOffset / DWORD_SIZE_IN_BYTES;
    uint MsfOffset;

    if( msfNumber == SCHEDULER_MSF_INITIAL )
    {
        MsfOffset = MEDIA_STATE_FLUSH_INITIAL_INTERFACE_DESCRIPTOR_OFFSET;
    }
    else if ( msfNumber == SCHEDULER_MSF_SECOND )
    {
        MsfOffset = MEDIA_STATE_FLUSH_INTERFACE_DESCRIPTOR_OFFSET;
    }
    else
    {
        return -1;
    }
    patchDword( &( secondaryBatchBuffer[ CmdPacketStart + MsfOffset ] ), 0, 5, interfaceDescriptorOffset );

    return 0;
}

#if defined WA_LRI_COMMANDS_EXIST
void PatchMiLoadRegisterImm(
    uint secondLevelBatchOffset,
    __global uint* secondaryBatchBuffer,
    uint enqueueOffset,
    uint registerAddress,
    uint value )
{
    //SlbOffset is expressed in bytes and for cmd it is needed to convert it to dwords
    uint CmdPacketStart = secondLevelBatchOffset / DWORD_SIZE_IN_BYTES;

    secondaryBatchBuffer[ CmdPacketStart + enqueueOffset ] = OCLRT_LOAD_REGISTER_IMM_CMD;
    patchDword( &( secondaryBatchBuffer[ CmdPacketStart + enqueueOffset + IMM_LOAD_REGISTER_ADDRESS_DWORD_OFFSET ] ), 2, 22, registerAddress >> 2 );
    secondaryBatchBuffer[ CmdPacketStart + enqueueOffset + IMM_LOAD_REGISTER_VALUE_DWORD_OFFSET ] = value;
}

void AddMiLoadRegisterImm(
    __global uint* secondaryBatchBuffer,
    __private uint* dwordOffset,
    uint value )
{
    secondaryBatchBuffer[ *dwordOffset ] = OCLRT_LOAD_REGISTER_IMM_CMD;
    ( *dwordOffset )++;
    secondaryBatchBuffer[ *dwordOffset ] = 0;
    patchDword( &( secondaryBatchBuffer[ *dwordOffset ] ), 2, 22, CTXT_PREMP_DBG_ADDRESS_VALUE >> 2 );
    ( *dwordOffset )++;
    secondaryBatchBuffer[ *dwordOffset ] = value; //CTXT_PREMP_ON_MI_ARB_CHECK_ONLY or CTXT_PREMP_DEFAULT_VALUE
    ( *dwordOffset )++;
}

void SetDisablePreemptionRegister(
    uint secondLevelBatchOffset,
    __global uint* secondaryBatchBuffer )
{
    PatchMiLoadRegisterImm( secondLevelBatchOffset,
                            secondaryBatchBuffer,
                            IMM_LOAD_REGISTER_FOR_DISABLE_PREEMPTION_OFFSET,
                            CTXT_PREMP_DBG_ADDRESS_VALUE,
                            CTXT_PREMP_ON_MI_ARB_CHECK_ONLY );
}

void SetEnablePreemptionRegister(
    uint secondLevelBatchOffset,
    __global uint* secondaryBatchBuffer )
{
    PatchMiLoadRegisterImm( secondLevelBatchOffset,
                            secondaryBatchBuffer,
                            IMM_LOAD_REGISTER_FOR_ENABLE_PREEMPTION_OFFSET,
                            CTXT_PREMP_DBG_ADDRESS_VALUE,
                            CTXT_PREMP_DEFAULT_VALUE );
}

void NoopPreemptionCommand(
    uint secondLevelBatchOffset,
    uint cmdOffset,
    __global uint* secondaryBatchBuffer )
{
    uint CmdPacketStart = cmdOffset + secondLevelBatchOffset / DWORD_SIZE_IN_BYTES;
    for( int i = 0; i < OCLRT_IMM_LOAD_REGISTER_CMD_DEVICE_CMD_DWORD_OFFSET; i++ )
    {
        secondaryBatchBuffer[ CmdPacketStart + i ] = 0;
    }
}
#endif //WA_LRI_COMMANDS_EXIST


//PQueue is needed for SLBOffset
void AddCmdsInSLBforScheduler20Parallel( uint slbOffset,
                                        __global IGIL_CommandQueue* pQueue,
                                        __global uint * secondaryBatchBuffer,
                                        __global char * dsh )
{
    EMULATION_ENTER_FUNCTION( );
#ifdef SCHEDULER_EMULATION
    uint3 StartPoint = { 0, 0, 0 };
    uint3 DimSize = { get_num_groups( 0 ), 1, 1 };
#else
    uint3 StartPoint = ( uint3 )( 0 );
    uint3 DimSize = ( uint3 )( get_num_groups( 0 ), 1, 1 );
#endif
    patchGpGpuWalker( slbOffset,
                      secondaryBatchBuffer,
                      0,
                      PARALLEL_SCHEDULER_COMPILATION_SIZE_20,
                      get_local_size(0),
                      DimSize,
                      StartPoint,
                      PARALLEL_SCHEDULER_HWTHREADS_IN_HW_GROUP20,
                      SIZEOF_3GRFS * PARALLEL_SCHEDULER_HWTHREADS_IN_HW_GROUP20 + pQueue->m_controls.m_SchedulerConstantBufferSize,
                      pQueue->m_controls.m_SchedulerDSHOffset );

    PatchMediaStateFlush( slbOffset, secondaryBatchBuffer, 0, SCHEDULER_MSF_INITIAL );
    PatchMediaStateFlush( slbOffset, secondaryBatchBuffer, 0, SCHEDULER_MSF_SECOND );

//When commands exists and scheduler does not require preemption off, noop the commands space
#if defined WA_LRI_COMMANDS_EXIST
#if defined WA_SCHEDULER_PREEMPTION
    if( pQueue->m_controls.m_EventTimestampAddress == 0u )
    {
        SetEnablePreemptionRegister( slbOffset, secondaryBatchBuffer );
        SetDisablePreemptionRegister( slbOffset, secondaryBatchBuffer );
    }
    else
    {
        NoopPreemptionCommand( slbOffset, IMM_LOAD_REGISTER_FOR_ENABLE_PREEMPTION_OFFSET, secondaryBatchBuffer );
        NoopPreemptionCommand( slbOffset, IMM_LOAD_REGISTER_FOR_DISABLE_PREEMPTION_OFFSET, secondaryBatchBuffer );
    }
#else
    //This is case, where LRI preemption is not required around scheduler WALKERs, but space for LRI commands exists, make sure they are nooped then
    NoopPreemptionCommand( SLBOffset, IMM_LOAD_REGISTER_FOR_ENABLE_PREEMPTION_OFFSET, secondaryBatchBuffer );
    NoopPreemptionCommand( SLBOffset, IMM_LOAD_REGISTER_FOR_DISABLE_PREEMPTION_OFFSET, secondaryBatchBuffer );
#endif //WA_SCHEDULER_PREEMPTION
#endif //WA_LRI_COMMANDS_EXIST
}

int generateLocalIDSParallel20(
    __global char* dsh,
    uint3 localSize,
    uint hwThreads,
    uint simdSize )
{
    uint it, currX, currY, currZ, FlattendID;

    uint Max = 1;
    if( simdSize == 32 )
    {
        Max = 2;
    }

    //Update full GRFs, each WI generate ID for one work item in x,y and z
    //in case we generate SIMD8 payload using 16 wi , idle half of them
    if( get_local_id( 0 ) < simdSize )
    {
        for( it = 0; it < hwThreads; it++ )
        {
            for( uint multip = 0; multip < Max; multip++ )
            {
                //We are in simd 8, each wi process generation for 1 wi
                FlattendID = get_local_id( 0 ) + it * simdSize + 16 * ( multip );

                currX = FlattendID % localSize.x;
                currY = ( FlattendID / localSize.x ) % localSize.y;
                currZ = ( FlattendID / ( localSize.x * localSize.y ) );//not needed % localSize.z;

                *( __global ushort * )( dsh + get_local_id( 0 ) * 2 + it * GRF_SIZE * 3 * Max + multip * GRF_SIZE )                                   = ( ushort )currX;
                *( __global ushort * )( dsh + get_local_id( 0 ) * 2 + it * GRF_SIZE * 3 * Max + GRF_SIZE * Max + multip * GRF_SIZE )                  = ( ushort )currY;
                *( __global ushort * )( dsh + get_local_id( 0 ) * 2 + it * GRF_SIZE * 3 * Max + GRF_SIZE * Max + GRF_SIZE * Max + multip * GRF_SIZE ) = ( ushort )currZ;
            }
        }
    }
    return 0;
}

//Function generate local ids.
//SIMD16 version
int generateLocalIDSsimd16(
    __global char* dsh,
    uint3 localSize,
    uint hwThreads)
{
    typedef union
    {
        ushort16 vectors;
        ushort varray[ 16 ];
    }vectorUnion;

    __private vectorUnion LidX;
    __private vectorUnion LidY;
    __private vectorUnion LidZ;

    __private ushort currX = 0;
    __private ushort currY = 0;
    __private ushort currZ = 0;

    //Assuming full load of hw thread , remainder done separately
    for(uint it = 0; it < hwThreads; it++ )
    {
        //This will be unrolled by compiler
        for(uint x = 0; x < 16; x++ )
        {
            LidX.varray[ x ] = currX++;
            LidY.varray[ x ] = currY;
            LidZ.varray[ x ] = currZ;

            if( currX == localSize.x )
            {
                currX = 0;
                currY++;
            }

            if( currY == localSize.y )
            {
                currY = 0;
                currZ++;
            }
        }

        *( __global ushort16 * )( dsh + it * GRF_SIZE * 3 )                          = LidX.vectors;
        *( __global ushort16 * )( dsh + it * GRF_SIZE * 3 + GRF_SIZE )               = LidY.vectors;
        *( __global ushort16 * )( dsh + it * GRF_SIZE * 3 + GRF_SIZE + GRF_SIZE )    = LidZ.vectors;

    }

    return 0;
}

//Function generate local ids.
//SIMD8 version
int generateLocalIDSsimd8(
    __global char* dsh,
    uint3 localSize,
    uint hwThreads)
{
    typedef union
    {
        ushort8 vectors;
        ushort varray[ 8 ];
    }vectorUnion;

    __private vectorUnion LidX;
    __private vectorUnion LidY;
    __private vectorUnion LidZ;

    __private ushort currX = 0;
    __private ushort currY = 0;
    __private ushort currZ = 0;

    //Assuming full load of hw thread , remainder done separately
    for(uint it = 0; it < hwThreads; it++ )
    {
        //This will be unrolled by compiler
        for(uint x = 0; x < 8; x++ )
        {
            LidX.varray[ x ] = currX++;
            LidY.varray[ x ] = currY;
            LidZ.varray[ x ] = currZ;

            if( currX == localSize.x )
            {
                currX = 0;
                currY++;
            }

            if( currY == localSize.y )
            {
                currY = 0;
                currZ++;
            }
        }

        *( __global ushort8 * )( dsh + it * GRF_SIZE * 3 )                          = LidX.vectors;
        *( __global ushort8 * )( dsh + it * GRF_SIZE * 3 + GRF_SIZE )               = LidY.vectors;
        *( __global ushort8 * )( dsh + it * GRF_SIZE * 3 + GRF_SIZE + GRF_SIZE )    = LidZ.vectors;
    }

    return 0;
}

//Function patches a curbe parametr , this version of function supports only these curbe tokens that may appear only once
int PatchDSH1Token( int currentIndex, uint tokenType, __global IGIL_KernelCurbeParams* pKernelCurbeParams, __global char* pDsh,
                    uint value )
{
    EMULATION_ENTER_FUNCTION( );

    uint PatchOffset;
#if SCHEDULER_DEBUG_MODE
    //If we are here it means that mask is ok and there are at least 3 curbe tokens that needs to be patched, do it right away
    if( pKernelCurbeParams[ CurrentIndex ].m_parameterType != TokenType )
    {
        return -1;
    }
#endif
    PatchOffset = pKernelCurbeParams[ currentIndex ].m_patchOffset;
    *( __global uint * )( &pDsh[ PatchOffset ] ) = value;

    currentIndex++;
    return currentIndex;
}

int PatchLocalMemEntities( int currentIndex, uint tokenType, __global IGIL_KernelCurbeParams* pKernelCurbeParams, __global char* pDsh,
                           __global IGIL_CommandHeader* pCommand )
{
    uint PatchOffset;
#if SCHEDULER_DEBUG_MODE
    //If we are here it means that mask is ok and there are at least 3 curbe tokens that needs to be patched, do it right away
    if( pKernelCurbeParams[ CurrentIndex ].m_parameterType != TokenType )
    {
        return -1;
    }
#endif
    //First patch is with 0
    PatchOffset  = pKernelCurbeParams[ currentIndex ].m_patchOffset;

    //SUM_OF_LOCAL_MEMORY_KERNEL_ARGS can be a 4 or 8 byte patch
    if( pKernelCurbeParams[currentIndex].m_parameterSize == sizeof( ulong ) )
    {
        *( __global ulong * )( &pDsh[PatchOffset] ) = 0;
    }
    else
    {
        *( __global uint * )( &pDsh[ PatchOffset ] ) = 0;
    }


    currentIndex++;
    uint Alignement;
    uint iter = 0;
    uint CurrentSum = 0;
    uint CurrentValue;
    //For each global captured there will be uint with index and ulong with address.
    uint GlobalPointersSize = ( pCommand->m_numGlobalCapturedBuffer * ( sizeof( ulong ) +  sizeof( uint ) ) ) / sizeof( uint );

    __global uint* pLocalMemSizes = &pCommand->m_data[ pCommand->m_numDependencies + pCommand->m_numScalarArguments + GlobalPointersSize ];
    //Check if there is second surface
    while( pKernelCurbeParams[ currentIndex ].m_parameterType == tokenType )
    {
        PatchOffset  = pKernelCurbeParams[ currentIndex ].m_patchOffset;

        //Value needs to be aligned to the value stored in sourceoffset
        Alignement = OCLRT_MAX( DWORD_SIZE_IN_BYTES, pKernelCurbeParams[ currentIndex ].m_sourceOffset );

        CurrentValue = pLocalMemSizes[ iter ];
        CurrentValue = OCLRT_ALIGN( CurrentValue, Alignement );
        CurrentSum   += CurrentValue;

        //SUM_OF_LOCAL_MEMORY_KERNEL_ARGS can be a 4 or 8 byte patch
        if( pKernelCurbeParams[currentIndex].m_parameterSize == sizeof( ulong ) )
        {
            *( __global ulong * )( &pDsh[PatchOffset] ) = ( ulong )CurrentSum;
        }
        else
        {
            *( __global uint * )( &pDsh[ PatchOffset ] ) = CurrentSum;
        }

        currentIndex++;
        iter++;
    }
    return currentIndex;
}

//Function patches a curbe parametr , this version of function supports only these curbe tokens that may appear only once
int PatchDSH1TokenParallel20( int currentIndex, uint tokenType, __global IGIL_KernelCurbeParams* pKernelCurbeParams, __global char* pDsh,
                              uint value )
{
    EMULATION_ENTER_FUNCTION( );

    uint PatchOffset;
    if( get_local_id( 0 ) == PARALLEL_SCHEDULER_COMPILATION_SIZE_20 )
    {
        PatchOffset  = pKernelCurbeParams[ currentIndex ].m_patchOffset;
        *( __global uint * ) ( &pDsh[ PatchOffset ] ) = value;
    }
    currentIndex++;
    return currentIndex;
}

//Function patches a curbe parametr, this version of function works on 3d curbe tokens
//It assumes that at least 3 tokens exists, then checks if 3 additional patches are needed
int PatchDSH6TokensParallel20( int currentIndex, uint tokenType, __global IGIL_KernelCurbeParams* pKernelCurbeParams, __global char* pDsh,
                               uint value1, uint value2, uint value3 )
{
    EMULATION_ENTER_FUNCTION( );

    uint PatchOffset, SourceOffset;
    uint WorkingOffset;
    uint ShiftSize;

    //Check if we patch 3 or 6 curbe tokens
    if( pKernelCurbeParams[ currentIndex + 3 ].m_parameterType == tokenType )
    {
        ShiftSize = 6;
    }
    else
    {
        ShiftSize = 3;
    }

    if( get_local_id( 0 ) < PARALLEL_SCHEDULER_COMPILATION_SIZE_20 + ShiftSize )
    {
        WorkingOffset = currentIndex + get_local_id( 0 ) - PARALLEL_SCHEDULER_COMPILATION_SIZE_20;
        PatchOffset   = pKernelCurbeParams[ WorkingOffset ].m_patchOffset;
        SourceOffset  = pKernelCurbeParams[ WorkingOffset ].m_sourceOffset;

        if( SourceOffset == 0 )
        {
            *( __global uint * )( &pDsh[ PatchOffset ] ) = value1;
        }
        else if( SourceOffset == 4 )
        {
            *( __global uint * )( &pDsh[ PatchOffset ] ) = value2;
        }
        else if( SourceOffset == 8 )
        {
            *( __global uint * )( &pDsh[ PatchOffset ] ) = value3;
        }
    }

    currentIndex += ShiftSize;
    return currentIndex;
}

int PatchLocalWorkSizes( int currentIndex, uint tokenType, __global IGIL_KernelCurbeParams* pKernelCurbeParams, __global char* pDsh,
                         uint enqLocalX, uint enqLocalY, uint enqLocalZ, uint cutLocalX, uint cutLocalY, uint cutLocalZ )
{
    EMULATION_ENTER_FUNCTION( );

    uint PatchOffset, SourceOffset;

    //Tokens are sorted by m_sourceOffset, it means that first 3 keys are always used to compute global_id and are always present
    for( uint it = 0; it < 3; it++ )
    {
        PatchOffset  = pKernelCurbeParams[ currentIndex ].m_patchOffset;
        SourceOffset = pKernelCurbeParams[ currentIndex ].m_sourceOffset;

        if( SourceOffset == 0 )
        {
            *( __global uint * )( &pDsh[ PatchOffset ] ) = enqLocalX;
        }
        else if( SourceOffset == 4 )
        {
            *( __global uint * )( &pDsh[ PatchOffset ] ) = enqLocalY;
        }
        else if( SourceOffset == 8 )
        {
            *( __global uint * )( &pDsh[ PatchOffset ] ) = enqLocalZ;
        }
        currentIndex++;
    }
    //If there are 3 more tokens, it means that get_local_size is used within a kernel, to deal with it patch with the second set of variables
    if( pKernelCurbeParams[ currentIndex ].m_parameterType == tokenType )
    {
        for( uint it = 0; it < 3; it++ )
        {
            PatchOffset  = pKernelCurbeParams[ currentIndex ].m_patchOffset;
            SourceOffset = pKernelCurbeParams[ currentIndex ].m_sourceOffset;

            if( SourceOffset == 0 )
            {
                *( __global uint * )( &pDsh[ PatchOffset ] ) = cutLocalX;
            }
            else if( SourceOffset == 4 )
            {
                *( __global uint * )( &pDsh[ PatchOffset ] ) = cutLocalY;
            }
            else if( SourceOffset == 8 )
            {
                *( __global uint * )( &pDsh[ PatchOffset ] ) = cutLocalZ;
            }
            currentIndex++;
        }
    }
    return currentIndex;
}

//Function patches a curbe parametr, this version of function works on 3d curbe tokens
//It assumes that at least 3 tokens exists, then checks if 3 additional patches are needed
int PatchLocalWorkSizesParallel( int currentIndex, uint tokenType, __global IGIL_KernelCurbeParams* pKernelCurbeParams, __global char* pDsh,
                                 uint enqLocalX, uint enqLocalY, uint enqLocalZ, uint cutLocalX, uint cutLocalY, uint cutLocalZ )
{
    EMULATION_ENTER_FUNCTION( );

    uint ShiftSize;

    //Check if we patch 3 or 6 curbe tokens
    if( pKernelCurbeParams[ currentIndex + 3 ].m_parameterType == tokenType )
    {
        ShiftSize = 6;
    }
    else
    {
        ShiftSize = 3;
    }

    //Use single threaded version
    if( get_local_id( 0 ) == PARALLEL_SCHEDULER_COMPILATION_SIZE_20 )
    {
        PatchLocalWorkSizes( currentIndex, SCHEDULER_DATA_PARAMETER_LOCAL_WORK_SIZE, pKernelCurbeParams, pDsh, enqLocalX, enqLocalY, enqLocalZ, cutLocalX, cutLocalY, cutLocalZ );
    }

    currentIndex += ShiftSize;
    return currentIndex;
}

//Function patches a curbe parametr, this version of function works on 3d curbe tokens
//It assumes that at least 3 tokens exists, then checks if 3 additional patches are needed
int PatchDSH6Tokens( int currentIndex, uint tokenType, __global IGIL_KernelCurbeParams* pKernelCurbeParams, __global char* pDsh,
                     uint value1, uint value2, uint value3 )
{
    EMULATION_ENTER_FUNCTION( );

    uint PatchOffset, SourceOffset;
#if SCHEDULER_DEBUG_MODE
    //If we are here it means that mask is ok and there are at least 3 curbe tokens that needs to be patched, do it right away
    if( pKernelCurbeParams[ CurrentIndex ].m_parameterType != TokenType )
    {
        return -1;
    }
#endif
    for( uint it = 0; it < 3; it++ )
    {
        PatchOffset  = pKernelCurbeParams[ currentIndex ].m_patchOffset;
        SourceOffset = pKernelCurbeParams[ currentIndex ].m_sourceOffset;

        if( SourceOffset == 0 )
        {
            *( __global uint * )( &pDsh[ PatchOffset ] ) = value1;
        }
        else if( SourceOffset == 4 )
        {
            *( __global uint * )( &pDsh[ PatchOffset ] ) = value2;
        }
        else if( SourceOffset == 8 )
        {
            *( __global uint * )( &pDsh[ PatchOffset ] ) = value3;
        }
        currentIndex++;
    }
    //Check if there are 3 more.
    if( pKernelCurbeParams[ currentIndex ].m_parameterType == tokenType )
    {
        for( uint it = 0; it < 3; it++ )
        {
            PatchOffset  = pKernelCurbeParams[ currentIndex ].m_patchOffset;
            SourceOffset = pKernelCurbeParams[ currentIndex ].m_sourceOffset;

            if( SourceOffset == 0 )
            {
                *( __global uint * )( &pDsh[ PatchOffset ] ) = value1;
            }
            else if( SourceOffset == 4 )
            {
                *( __global uint * )( &pDsh[ PatchOffset ] ) = value2;
            }
            else if( SourceOffset == 8 )
            {
                *( __global uint * )( &pDsh[ PatchOffset ] ) = value3;
            }
            currentIndex++;
        }
    }

    return currentIndex;
}
//Common code

inline __global char* GetPtrToCurbeData( uint offset, __global IGIL_KernelDataHeader * pKernelReflection )
{
    __global char * pRawKernelReflection = ( __global char * )pKernelReflection;
    return ( pRawKernelReflection + offset );
}

__global char* GetPtrToKernelReflectionOffset( uint offset, __global IGIL_KernelDataHeader * pKernelReflection )
{
    __global char * pRawKernelReflection = ( __global char * )pKernelReflection;
    return ( pRawKernelReflection + offset );
}

void InitWalkerDataParallel( __local IGIL_WalkerEnumeration* pWalkerEnumData,
                             uint workDim,
                             uint* pWalkerCount,
                             uint3 edgeArray,
                             uint3 globalDim,
                             uint3 globalSizes,
                             uint3 localSizes )
{
    EMULATION_ENTER_FUNCTION( );

    pWalkerEnumData->TotalDimSize.x = globalDim.x;
    pWalkerEnumData->TotalDimSize.y = globalDim.y;
    pWalkerEnumData->TotalDimSize.z = globalDim.z;

    pWalkerEnumData->WalkerArray[ 0 ].ActualLocalSize.x = localSizes.x;
    pWalkerEnumData->WalkerArray[ 0 ].WalkerStartPoint.x = 0;
    pWalkerEnumData->WalkerArray[ 0 ].WalkerDimSize.x = globalDim.x;

    pWalkerEnumData->WalkerArray[ 0 ].ActualLocalSize.y = localSizes.y;
    pWalkerEnumData->WalkerArray[ 0 ].WalkerStartPoint.y = 0;
    pWalkerEnumData->WalkerArray[ 0 ].WalkerDimSize.y = globalDim.y;

    pWalkerEnumData->WalkerArray[ 0 ].ActualLocalSize.z = localSizes.z;
    pWalkerEnumData->WalkerArray[ 0 ].WalkerStartPoint.z = 0;
    pWalkerEnumData->WalkerArray[ 0 ].WalkerDimSize.z = globalDim.z;

    uint WalkerCount = 1;

    if( edgeArray.x != 0 )
    {
        pWalkerEnumData->TotalDimSize.x++;

        pWalkerEnumData->WalkerArray[ 1 ].ActualLocalSize.x = edgeArray.x;
        pWalkerEnumData->WalkerArray[ 1 ].WalkerStartPoint.x = globalDim.x;
        pWalkerEnumData->WalkerArray[ 1 ].WalkerDimSize.x = pWalkerEnumData->TotalDimSize.x;

        pWalkerEnumData->WalkerArray[ 1 ].ActualLocalSize.y = localSizes.y;
        pWalkerEnumData->WalkerArray[ 1 ].WalkerStartPoint.y = 0;
        pWalkerEnumData->WalkerArray[ 1 ].WalkerDimSize.y = globalDim.y;

        pWalkerEnumData->WalkerArray[ 1 ].ActualLocalSize.z = localSizes.z;
        pWalkerEnumData->WalkerArray[ 1 ].WalkerStartPoint.z = 0;
        pWalkerEnumData->WalkerArray[ 1 ].WalkerDimSize.z = globalDim.z;

        WalkerCount++;
    }

    if( workDim > 1 )
    {
        if( edgeArray.y != 0 )
        {
            pWalkerEnumData->TotalDimSize.y++;

            pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.x = localSizes.x;
            pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.x = 0;
            pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.x = globalDim.x;

            pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.y = edgeArray.y;
            pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.y = globalDim.y;
            pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.y = pWalkerEnumData->TotalDimSize.y;

            pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.z = localSizes.z;
            pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.z = 0;
            pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.z = globalDim.z;

            WalkerCount++;
        }

        if( ( edgeArray.x != 0 ) & ( edgeArray.y != 0 ) )
        {
            pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.x = edgeArray.x;
            pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.x = globalDim.x;
            pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.x = pWalkerEnumData->TotalDimSize.x;

            pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.y = edgeArray.y;
            pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.y = globalDim.y;
            pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.y = pWalkerEnumData->TotalDimSize.y;

            pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.z = localSizes.z;
            pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.z = 0;
            pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.z = globalDim.z;

            WalkerCount++;
        }
        if( workDim > 2 )
        {
            if( edgeArray.z != 0 )
            {
                pWalkerEnumData->TotalDimSize.z++;

                pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.x = localSizes.x;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.x = 0;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.x = globalDim.x;

                pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.y = localSizes.y;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.y = 0;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.y = globalDim.y;

                pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.z = edgeArray.z;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.z = globalDim.z;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.z = pWalkerEnumData->TotalDimSize.z;

                WalkerCount++;
            }
            if( ( edgeArray.x != 0 ) & ( edgeArray.z != 0 ) )
            {
                pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.x = edgeArray.x;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.x = globalDim.x;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.x = pWalkerEnumData->TotalDimSize.x;

                pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.y = localSizes.y;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.y = 0;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.y = globalDim.y;

                pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.z = edgeArray.z;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.z = globalDim.z;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.z = pWalkerEnumData->TotalDimSize.z;

                WalkerCount++;
            }

            if( ( edgeArray.y != 0 ) & ( edgeArray.z != 0 ) )
            {
                pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.x = localSizes.x;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.x = 0;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.x = globalDim.x;

                pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.y = edgeArray.y;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.y = globalDim.y;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.y = pWalkerEnumData->TotalDimSize.y;

                pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.z = edgeArray.z;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.z = globalDim.z;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.z = pWalkerEnumData->TotalDimSize.z;

                WalkerCount++;
            }
            if( ( edgeArray.x != 0 ) & ( edgeArray.y != 0 ) & ( edgeArray.z != 0 ) )
            {
                pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.x = edgeArray.x;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.x = globalDim.x;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.x = pWalkerEnumData->TotalDimSize.x;

                pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.y = edgeArray.y;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.y = globalDim.y;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.y = pWalkerEnumData->TotalDimSize.y;

                pWalkerEnumData->WalkerArray[ WalkerCount ].ActualLocalSize.z = edgeArray.z;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerStartPoint.z = globalDim.z;
                pWalkerEnumData->WalkerArray[ WalkerCount ].WalkerDimSize.z = pWalkerEnumData->TotalDimSize.z;

                WalkerCount++;
            }
        }
    }

    *pWalkerCount = WalkerCount;
}

//Compute number of Walkers needed for this command packet, this function assumes that command packet is initialized
inline int GetWalkerCount( __global IGIL_CommandHeader* pCommand )
{
    int WalkerCount = 1;
    for( uint dim = 0; ( dim < pCommand->m_range.m_dispatchDimensions ); dim++ )
    {
        if( ( pCommand->m_range.m_globalWorkSize[ dim ] % pCommand->m_range.m_localWorkSize[ dim ] ) != 0 )
        {
            WalkerCount *= 2;
        }
    }
    return WalkerCount;
}

//This function intializes command packet, checks for null case, sets proper LWS sizes and return WalkerCount needed for this packet
inline void InitializeCommandPacket( __global IGIL_CommandHeader* pCommand )
{
    EMULATION_ENTER_FUNCTION( );

    //Check for NULL case
    if( pCommand->m_range.m_localWorkSize[ 0 ] == 0 )
    {
        //If null case detected use 16 x 1 x 1 lws
        if( pCommand->m_range.m_globalWorkSize[ 0 ] >= 16 )
        {
            pCommand->m_range.m_localWorkSize[ 0 ] = 16;
        }
        else
        {
            pCommand->m_range.m_localWorkSize[ 0 ] = pCommand->m_range.m_globalWorkSize[ 0 ];
        }

        pCommand->m_range.m_localWorkSize[ 1 ] = 1;
        pCommand->m_range.m_localWorkSize[ 2 ] = 1;
    }

}

//Patches the address for pipe  control
void PatchPipeControlProfilingAddres( __global uint* secondaryBatchBuffer, uint slBoffset, ulong address, uint pipeControlOffset )
{
    EMULATION_ENTER_FUNCTION( );

    //SlbOffset is expressed in bytes and for cmd it is needed to convert it to dwords
    uint PostSyncDwordOffset = ( slBoffset / DWORD_SIZE_IN_BYTES ) + pipeControlOffset + PIPE_CONTROL_POST_SYNC_DWORD;
    uint DwordOffset         = ( slBoffset / DWORD_SIZE_IN_BYTES ) + pipeControlOffset + PIPE_CONTROL_ADDRESS_FIELD_DWORD;
    //Patch P_C event timestamp address in SLB in 3rd and 4th dword
    secondaryBatchBuffer[ DwordOffset ] = 0;
    patchDword( &secondaryBatchBuffer[ DwordOffset ], PIPE_CONTROL_GRAPHICS_ADDRESS_START_BIT, PIPE_CONTROL_GRAPHICS_ADDRESS_END_BIT, ( uint )( address >> PIPE_CONTROL_GRAPHICS_ADDRESS_START_BIT ) );
    DwordOffset++;
    secondaryBatchBuffer[ DwordOffset ] = 0;
    patchDword( &secondaryBatchBuffer[ DwordOffset ], PIPE_CONTROL_GRAPHICS_ADDRESS_HIGH_START_BIT, PIPE_CONTROL_GRAPHICS_ADDRESS_HIGH_END_BIT, ( address >> 32 ) );

    //Patch Timestamp bit
    patchDword( &secondaryBatchBuffer[ PostSyncDwordOffset ], PIPE_CONTROL_POST_SYNC_START_BIT, PIPE_CONTROL_POST_SYNC_END_BIT, PIPE_CONTROL_GENERATE_TIME_STAMP );
}

void DisablePostSyncBitInPipeControl( __global uint* secondaryBatchBuffer, uint slBoffset, uint pipeControlOffset )
{
    //SlbOffset is expressed in bytes and for cmd it is needed to convert it to dwords
    uint PostSyncDwordOffset = ( slBoffset / DWORD_SIZE_IN_BYTES ) + pipeControlOffset + PIPE_CONTROL_POST_SYNC_DWORD;
    //Patch P_C event timestamp address in SLB in 3rd and 4th dword
    patchDword( &secondaryBatchBuffer[ PostSyncDwordOffset ], PIPE_CONTROL_POST_SYNC_START_BIT, PIPE_CONTROL_POST_SYNC_END_BIT, PIPE_CONTROL_NO_POSTSYNC_OPERATION );
}


int PatchDSH( __global IGIL_CommandQueue* pQueue,
              __global IGIL_KernelDataHeader * pKernelReflection,
              __global char* dsh,
              uint blockId,
              __global IGIL_CommandHeader* pCommandHeader,
              __global uint* secondaryBatchBuffer,
              uint dshOffset,
              uint intefaceDescriptorOffset,
              __local IGIL_WalkerEnumeration* pWalkerMain,
              uint walkerPos )
{
    EMULATION_ENTER_FUNCTION( );

    __global IGIL_KernelAddressData* pKernelAddressData = IGIL_GetKernelAddressData( pKernelReflection, blockId );
    __global IGIL_KernelData* pBlockData                = IGIL_GetKernelData( pKernelReflection, blockId );
    ulong PatchMask                                     = pBlockData->m_PatchTokensMask;
    uint CurrentIndex                                   = 0;

    __global char* pDsh                                 = ( __global char* )&dsh[ dshOffset ];
    __global IGIL_KernelCurbeParams* pKernelCurbeParams = ( __global IGIL_KernelCurbeParams* )&pBlockData->m_data;

    uint NumberOfDepencies                              = pCommandHeader->m_numDependencies;
    uint PatchOffset;
    __global char* pScalarData                          = ( __global char* )( &pCommandHeader->m_data[ NumberOfDepencies ] );
    __global char *pDshOnKRS                            = GetPtrToKernelReflectionOffset( pKernelAddressData->m_SamplerHeapOffset, pKernelReflection );

    uint SizeOfScalarsFromCurbe                         = 0;
    uint CurbeSize;

    uint TotalLocalSize;
    uint ThreadPayloadSize;
    uint NumberOfHWThreads;
    uint WorkDim;
    uint3 GlobalOffset;
    uint3 GlobalSizes;
    uint3 ActualLocalSize;

    GlobalOffset.x = ( uint )pCommandHeader->m_range.m_globalWorkOffset[ 0 ];
    GlobalOffset.y = ( uint )pCommandHeader->m_range.m_globalWorkOffset[ 1 ];
    GlobalOffset.z = ( uint )pCommandHeader->m_range.m_globalWorkOffset[ 2 ];

    GlobalSizes.x = ( uint )pCommandHeader->m_range.m_globalWorkSize[ 0 ];
    GlobalSizes.y = ( uint )pCommandHeader->m_range.m_globalWorkSize[ 1 ];
    GlobalSizes.z = ( uint )pCommandHeader->m_range.m_globalWorkSize[ 2 ];

    ActualLocalSize.x = pWalkerMain->WalkerArray[ walkerPos ].ActualLocalSize.x;
    ActualLocalSize.y = pWalkerMain->WalkerArray[ walkerPos ].ActualLocalSize.y;
    ActualLocalSize.z = pWalkerMain->WalkerArray[ walkerPos ].ActualLocalSize.z;

    WorkDim = pCommandHeader->m_range.m_dispatchDimensions;
    TotalLocalSize = ActualLocalSize.x * ActualLocalSize.y * ActualLocalSize.z;

    NumberOfHWThreads = TotalLocalSize / pBlockData->m_SIMDSize;
    if( TotalLocalSize % pBlockData->m_SIMDSize )
    {
        NumberOfHWThreads++;
    }

    ThreadPayloadSize = NumberOfHWThreads * 3 * GRF_SIZE;
    //Copy constant buffer to designated area on DSH.

    //pDshOnKRS seems to be in the global address space, not the private address space
    //Copy SamplerState and Constant Buffer at once
    IGILLOCAL_MEMCPY_GTOG( pDsh, pDshOnKRS, pBlockData->m_sizeOfConstantBuffer + pBlockData->m_SizeOfSamplerHeap );

    if( PatchMask & SCHEDULER_DATA_PARAMETER_KERNEL_ARGUMENT_MASK )
    {
        while( pKernelCurbeParams[ CurrentIndex ].m_parameterType == SCHEDULER_DATA_PARAMETER_KERNEL_ARGUMENT )
        {
            CurbeSize               = pKernelCurbeParams[ CurrentIndex ].m_parameterSize;
            SizeOfScalarsFromCurbe += CurbeSize;
            PatchOffset             = pKernelCurbeParams[ CurrentIndex ].m_patchOffset;

            //pScalarData is in the global address space, not the private address space
            IGILLOCAL_MEMCPY_GTOG( &pDsh[ PatchOffset ], pScalarData, CurbeSize );
            pScalarData            += CurbeSize;
            CurrentIndex++;
        }
#if SCHEDULER_DEBUG_MODE
        if( pCommandHeader->m_sizeOfScalarArguments != SizeOfScalarsFromCurbe )
        {
            pCommandHeader->m_commandState = SCHEDULER_CURBE_ARGUMENTS_SIZE_MISMATCH;
            return -1;
        }
#endif
    }

    if( PatchMask & SCHEDULER_DATA_PARAMETER_LOCAL_WORK_SIZE_MASK )
    {
        CurrentIndex     = PatchLocalWorkSizes( CurrentIndex,
                                                SCHEDULER_DATA_PARAMETER_LOCAL_WORK_SIZE,
                                                pKernelCurbeParams,
                                                pDsh,
                                                pWalkerMain->LocalWorkSize.x,
                                                pWalkerMain->LocalWorkSize.y,
                                                pWalkerMain->LocalWorkSize.z,
                                                ActualLocalSize.x,
                                                ActualLocalSize.y,
                                                ActualLocalSize.z );
#if SCHEDULER_DEBUG_MODE
        if( ( CurrentIndex == -1 ) || ( CurrentIndex >= pBlockData->m_numberOfCurbeParams ) )
        {
            pCommandHeader->m_commandState = SCHEDULER_CURBE_TOKEN_MISSED;
            return -1;
        }
#endif
    }

    if( PatchMask & SCHEDULER_DATA_PARAMETER_GLOBAL_WORK_SIZE_MASK )
    {
        CurrentIndex     = PatchDSH6Tokens( CurrentIndex,
                                            SCHEDULER_DATA_PARAMETER_GLOBAL_WORK_SIZE,
                                            pKernelCurbeParams,
                                            pDsh,
                                            GlobalSizes.x,
                                            GlobalSizes.y,
                                            GlobalSizes.z );
#if SCHEDULER_DEBUG_MODE
        if( ( CurrentIndex == -1 ) || ( CurrentIndex >= pBlockData->m_numberOfCurbeParams ) )
        {
            pCommandHeader->m_commandState = SCHEDULER_CURBE_TOKEN_MISSED;
            return -1;
        }
#endif
    }
    if( PatchMask & SCHEDULER_DATA_PARAMETER_NUM_WORK_GROUPS_MASK )
    {
        CurrentIndex     = PatchDSH6Tokens( CurrentIndex,
                                            SCHEDULER_DATA_PARAMETER_NUM_WORK_GROUPS,
                                            pKernelCurbeParams,
                                            pDsh,
                                            pWalkerMain->TotalDimSize.x,
                                            pWalkerMain->TotalDimSize.y,
                                            pWalkerMain->TotalDimSize.z );
#if SCHEDULER_DEBUG_MODE
        if( ( CurrentIndex == -1 ) || ( CurrentIndex >= pBlockData->m_numberOfCurbeParams ) )
        {
            pCommandHeader->m_commandState = SCHEDULER_CURBE_TOKEN_MISSED;
            return -1;
        }
#endif
    }
    if( PatchMask & SCHEDULER_DATA_PARAMETER_WORK_DIMENSIONS_MASK )
    {
        CurrentIndex     = PatchDSH1Token( CurrentIndex,
                                           SCHEDULER_DATA_PARAMETER_WORK_DIMENSIONS,
                                           pKernelCurbeParams,
                                           pDsh,
                                           WorkDim );
#if SCHEDULER_DEBUG_MODE
        if( ( CurrentIndex == -1 ) || ( CurrentIndex >= pBlockData->m_numberOfCurbeParams ) )
        {
            pCommandHeader->m_commandState = SCHEDULER_CURBE_TOKEN_MISSED;
            return -1;
        }
#endif
    }

    if( PatchMask & SCHEDULER_DATA_PARAMETER_SUM_OF_LOCAL_MEMORY_OBJECT_ARGUMENT_SIZES_MASK )
    {
        CurrentIndex     = PatchLocalMemEntities( CurrentIndex,
                                                  SCHEDULER_DATA_PARAMETER_SUM_OF_LOCAL_MEMORY_OBJECT_ARGUMENT_SIZES,
                                                  pKernelCurbeParams,
                                                  pDsh,
                                                  pCommandHeader );
#if SCHEDULER_DEBUG_MODE
        if( ( CurrentIndex == -1 ) || ( CurrentIndex >= pBlockData->m_numberOfCurbeParams ) )
        {
            pCommandHeader->m_commandState = SCHEDULER_CURBE_TOKEN_MISSED;
            return -1;
        }
#endif
    }

    if( PatchMask & SCHEDULER_DATA_PARAMETER_GLOBAL_WORK_OFFSET_MASK )
    {
        CurrentIndex     = PatchDSH6Tokens( CurrentIndex,
                                            SCHEDULER_DATA_PARAMETER_GLOBAL_WORK_OFFSET,
                                            pKernelCurbeParams,
                                            pDsh,
                                            GlobalOffset.x,
                                            GlobalOffset.y,
                                            GlobalOffset.z );
#if SCHEDULER_DEBUG_MODE
        if( ( CurrentIndex == -1 ) || ( CurrentIndex >= pBlockData->m_numberOfCurbeParams ) )
        {
            pCommandHeader->m_commandState = SCHEDULER_CURBE_TOKEN_MISSED;
            return -1;
        }
#endif
    }

    if( PatchMask & SCHEDULER_DATA_PARAMETER_NUM_HARDWARE_THREADS_MASK )
    {
        CurrentIndex     = PatchDSH1Token( CurrentIndex,
                                           SCHEDULER_DATA_PARAMETER_NUM_HARDWARE_THREADS,
                                           pKernelCurbeParams,
                                           pDsh,
                                           NumberOfHWThreads );
#if SCHEDULER_DEBUG_MODE
        if( ( CurrentIndex == -1 ) || ( CurrentIndex >= pBlockData->m_numberOfCurbeParams ) )
        {
            pCommandHeader->m_commandState = SCHEDULER_CURBE_TOKEN_MISSED;
            return -1;
        }
#endif
    }

    if( PatchMask & SCHEDULER_DATA_PARAMETER_PARENT_EVENT_MASK )
    {
        CurrentIndex     = PatchDSH1Token( CurrentIndex,
                                           SCHEDULER_DATA_PARAMETER_PARENT_EVENT,
                                           pKernelCurbeParams,
                                           pDsh,
                                           pCommandHeader->m_event );
#if SCHEDULER_DEBUG_MODE
        if( ( CurrentIndex == -1 ) || ( CurrentIndex >= pBlockData->m_numberOfCurbeParams ) )
        {
            pCommandHeader->m_commandState = SCHEDULER_CURBE_TOKEN_MISSED;
            return -1;
        }
#endif
    }

    if( PatchMask & SCHEDULER_DATA_PARAMETER_ENQUEUED_LOCAL_WORK_SIZE_MASK )
    {
        CurrentIndex     = PatchDSH6Tokens( CurrentIndex,
                                            SCHEDULER_DATA_PARAMETER_ENQUEUED_LOCAL_WORK_SIZE,
                                            pKernelCurbeParams,
                                            pDsh,
                                            pWalkerMain->LocalWorkSize.x,
                                            pWalkerMain->LocalWorkSize.y,
                                            pWalkerMain->LocalWorkSize.z );
#if SCHEDULER_DEBUG_MODE
        if( ( CurrentIndex == -1 ) || ( CurrentIndex >= pBlockData->m_numberOfCurbeParams ) )
        {
            pCommandHeader->m_commandState = SCHEDULER_CURBE_TOKEN_MISSED;
            return -1;
        }
#endif
    }
    if( PatchMask & SCHEDULER_DATA_PARAMETER_GLOBAL_POINTER )
    {
        if( pCommandHeader->m_numGlobalCapturedBuffer > 0 )
        {
            //Handle global pointers patching in stateless mode, info about layout in declaration of IGIL_CommandHeader
            __global    uint*  pGlobalIndexes = ( __global uint* ) ( &pCommandHeader->m_data[ NumberOfDepencies + pCommandHeader->m_numScalarArguments ] );
            __global    uint*  pGlobalPtrs    = ( __global uint* ) ( &pCommandHeader->m_data[ NumberOfDepencies + pCommandHeader->m_numScalarArguments + pCommandHeader->m_numGlobalCapturedBuffer ] );
            uint        StartIndex            = CurrentIndex;

            //Argument in command header are not in correct sequence, that's why proper key needs to be located
            for( uint glIdx = 0 ; glIdx < pCommandHeader->m_numGlobalCapturedBuffer; glIdx++)
            {
                //Reset CurrentIndex as we need to start from the beginning.
                CurrentIndex  = StartIndex;
                while( pKernelCurbeParams[ CurrentIndex ].m_parameterType == COMPILER_DATA_PARAMETER_GLOBAL_SURFACE )
                {
                    //Patch only if exact match occurs
                    if( pKernelCurbeParams[ CurrentIndex ].m_sourceOffset == *pGlobalIndexes )
                    {
                        PatchOffset             = pKernelCurbeParams[ CurrentIndex ].m_patchOffset;
                        //64 bit patching
                        if( pKernelCurbeParams[ CurrentIndex ].m_parameterSize == 8 )
                        {
                            __global uint* pDst = (__global uint *) &pDsh[PatchOffset];
                            pDst[ 0 ] = pGlobalPtrs[ 0 ];
                            pDst[ 1 ] = pGlobalPtrs[ 1 ];
                        }
                        else
                        {
                            __global uint* pDst = ( __global uint* ) &pDsh[ PatchOffset ];
                            *pDst               = ( uint ) *pGlobalPtrs;
                        }
                    }
                    CurrentIndex++;
                }
                pGlobalPtrs += 2;
                pGlobalIndexes++;
            }
        }
    }

    //Now generate local IDS
    if( pBlockData->m_SIMDSize == 8 )
    {
        generateLocalIDSsimd8( &pDsh[ pBlockData->m_sizeOfConstantBuffer ], ActualLocalSize, NumberOfHWThreads );
    }
    else
    {
        generateLocalIDSsimd16( &pDsh[ pBlockData->m_sizeOfConstantBuffer ], ActualLocalSize, NumberOfHWThreads );
    }

    uint TotalSLMSize = pCommandHeader->m_totalLocalSize + pBlockData->m_InilineSLMSize;

    //Update Interface Descriptor Data with SLM size  / number of HW threads.
    CopyAndPatchIDData( dsh, blockId, NumberOfHWThreads, TotalSLMSize, intefaceDescriptorOffset, pQueue->m_controls.m_StartBlockID );

    //Add WalkerStartSize
    patchGpGpuWalker( pQueue->m_controls.m_SecondLevelBatchOffset, secondaryBatchBuffer, intefaceDescriptorOffset, pBlockData->m_SIMDSize,
                      TotalLocalSize, pWalkerMain->WalkerArray[ walkerPos ].WalkerDimSize, pWalkerMain->WalkerArray[ walkerPos ].WalkerStartPoint,
                      NumberOfHWThreads, pBlockData->m_sizeOfConstantBuffer + ThreadPayloadSize, dshOffset );

    PatchMediaStateFlush( pQueue->m_controls.m_SecondLevelBatchOffset, secondaryBatchBuffer, intefaceDescriptorOffset, SCHEDULER_MSF_INITIAL );
    PatchMediaStateFlush( pQueue->m_controls.m_SecondLevelBatchOffset, secondaryBatchBuffer, intefaceDescriptorOffset, SCHEDULER_MSF_SECOND );

    return 0;
}

//Returns: isSRGB(ChannelOrder) ? ChannelOrder : 0;
inline uint GetSRGBChannelOrder( uint channelOrder )
{
    const uint AsSrgb = channelOrder - CL_sRGB;
    const uint NumSrgbFormats = CL_sBGRA - CL_sRGB;
    if( AsSrgb < NumSrgbFormats )
      return channelOrder;
    else
      return 0;
}

void PatchDSHParallelWithDynamicDSH20( uint slbOffsetBase,
                                      uint dshOffsetBase,
                                      uint intefaceDescriptorOffsetBase,
                                      __global IGIL_KernelDataHeader * pKernelReflection,
                                      __global char* dsh,
                                      uint blockId,
                                      __global IGIL_CommandHeader* pCommandHeader,
                                      __global uint* secondaryBatchBuffer,
                                      __global IGIL_CommandQueue* pQueue,
                                      __global IGIL_EventPool* eventsPool,
                                      __global char* ssh,
                                      uint btOffset,
                                      __local IGIL_WalkerEnumeration* pWalkerEnum,
                                      __local uint* objectIds
#ifdef ENABLE_DEBUG_BUFFER
                                      , __global DebugDataBuffer* DebugQueue
#endif
                                      )
{
    EMULATION_ENTER_FUNCTION( );

    __global IGIL_KernelAddressData* pKernelAddressData = IGIL_GetKernelAddressData( pKernelReflection, blockId );
    __global IGIL_KernelData* pBlockData                = IGIL_GetKernelData( pKernelReflection, blockId );
    ulong PatchMask                                     = pBlockData->m_PatchTokensMask;
    uint CurrentIndex                                   = 0;

    __global IGIL_KernelCurbeParams* pKernelCurbeParams = ( __global IGIL_KernelCurbeParams* )&pBlockData->m_data;

    uint NumberOfDepencies                              = pCommandHeader->m_numDependencies;
    uint PatchOffset;

    uint CurbeSize;

    //Get pointer to the Sampler State
    __global char *pDshOnKRS                            = GetPtrToKernelReflectionOffset( pKernelAddressData->m_SamplerHeapOffset, pKernelReflection );

    uint WalkerCount                                    = GetWalkerCount( pCommandHeader );
    __global char *pKernelReflectionChar                = ( __global char * ) pKernelReflection;
    __global IGIL_KernelCurbeParams* pSSHdata           = ( __global  IGIL_KernelCurbeParams* )&pKernelReflectionChar[ pKernelAddressData->m_SSHTokensOffset ];

    //WALKER variables that will be propagated to SLB
    uint3 LocalSizes;
    uint3 GlobalSizes;
    uint3 GlobalOffset;
    uint3 EdgeArray;
    uint3 XYZDim;

    //X is always there
    GlobalOffset.x  = ( uint )pCommandHeader->m_range.m_globalWorkOffset[ 0 ];
    GlobalSizes.x   = ( uint )pCommandHeader->m_range.m_globalWorkSize[ 0 ];
    LocalSizes.x    = ( uint )pCommandHeader->m_range.m_localWorkSize[ 0 ];
    EdgeArray.x     = GlobalSizes.x % LocalSizes.x;
    uint WorkDim    = pCommandHeader->m_range.m_dispatchDimensions;
    XYZDim.x        = GlobalSizes.x / LocalSizes.x;

    if( WorkDim > 1 )
    {
        GlobalOffset.y  = ( uint )pCommandHeader->m_range.m_globalWorkOffset[ 1 ];
        GlobalSizes.y   = ( uint )pCommandHeader->m_range.m_globalWorkSize[ 1 ];
        LocalSizes.y    = ( uint )pCommandHeader->m_range.m_localWorkSize[ 1 ];
        EdgeArray.y     = GlobalSizes.y % LocalSizes.y;
        XYZDim.y        = GlobalSizes.y / LocalSizes.y;

        if( WorkDim > 2 )
        {
            GlobalOffset.z  = ( uint )pCommandHeader->m_range.m_globalWorkOffset[ 2 ];
            GlobalSizes.z   = ( uint )pCommandHeader->m_range.m_globalWorkSize[ 2 ];
            LocalSizes.z    = ( uint )pCommandHeader->m_range.m_localWorkSize[ 2 ];
            XYZDim.z        = GlobalSizes.z / LocalSizes.z;
            EdgeArray.z     = GlobalSizes.z % LocalSizes.z;
        }
        else
        {
            GlobalOffset.z  = 0;
            GlobalSizes.z   = 1;
            LocalSizes.z    = 1;
            EdgeArray.z     = 0;
            XYZDim.z        = 1;
        }
    }
    else
    {
        GlobalOffset.y  = 0;
        GlobalOffset.z  = 0;
        GlobalSizes.y   = 1;
        GlobalSizes.z   = 1;
        LocalSizes.y    = 1;
        LocalSizes.z    = 1;
        EdgeArray.z     = 0;
        EdgeArray.y     = 0;
        XYZDim.z        = 1;
        XYZDim.y        = 1;
    }

    if( get_local_id( 0 ) == 0 )
    {
        InitWalkerDataParallel( pWalkerEnum, WorkDim, &WalkerCount, EdgeArray, XYZDim, GlobalSizes, LocalSizes );
    }

    uint SLBOffset                  = slbOffsetBase;
    uint DshOffset                  = dshOffsetBase;
    uint IntefaceDescriptorOffset   = intefaceDescriptorOffsetBase;
    __global uint* pArgumentIds     = NULL;
    __global uint* pObjectIds       = NULL;
    __global char* pLocalIdsOnDSH   = NULL;

    uint SamplerHeapSize = pBlockData->m_SizeOfSamplerHeap;

    //Object ID is in fact surface state offset for parent in case of surfaces using SSH, copy SSH from parent to child.
    //Copy binding table state of this kernel to allocated place on ssh
    GLOBAL_MEMCPY( &ssh[ btOffset ], &ssh[ pKernelAddressData->m_BTSoffset ] , pKernelAddressData->m_BTSize );

    for( uint WalkerID = 0; WalkerID < WalkerCount; WalkerID++ )
    {
        //Update the offsets
        if( WalkerID > 0 )
        {
            SLBOffset += SECOND_LEVEL_BUFFER_SPACE_FOR_EACH_ENQUEUE;
            SLBOffset %= MAX_SLB_OFFSET;
            IntefaceDescriptorOffset++;
            DshOffset += MAX_DSH_SIZE_PER_ENQUEUE;
        }

        __global char* pDsh = ( __global char* )&dsh[ DshOffset ];
        pLocalIdsOnDSH      = &pDsh[ pBlockData->m_sizeOfConstantBuffer + SamplerHeapSize ];

        //Copy Sampler State  and constant buffer on all threads
        GLOBAL_MEMCPY( pDsh, pDshOnKRS, pBlockData->m_sizeOfConstantBuffer + SamplerHeapSize );

        barrier( CLK_GLOBAL_MEM_FENCE );

        //Update BorderColorPointer on all threads
        if( pBlockData->m_numberOfSamplerStates )
        {
            uint SamplerId = get_local_id( 0 );
            __global uint* pSamplerState;
            while( SamplerId < pBlockData->m_numberOfSamplerStates )
            {
                pSamplerState = ( __global uint* )&dsh[ DshOffset + pBlockData->m_SamplerStateArrayOffsetOnDSH + SamplerId * OCLRT_SIZEOF_SAMPLER_STATE ];
                uint PatchValue = DshOffset >> 5;
                patchDword( &pSamplerState[ SAMPLER_STATE_DESCRIPTOR_BORDER_COLOR_POINTER_DWORD ], 5, 31, PatchValue );
                SamplerId += PARALLEL_SCHEDULER_COMPILATION_SIZE_20 * PARALLEL_SCHEDULER_HWTHREADS_IN_HW_GROUP20;
            }
        }

        //Setup SSH if needed, do it only for first Walker as all Walkers will re-use the same binding table layout.
        if( ( pCommandHeader->m_numGlobalArguments > 0 ) & ( WalkerID == 0 ) )
        {
            //Global arguments are after scalars, global pointers slm sizes and events
            uint offset                 = pCommandHeader->m_numDependencies + pCommandHeader->m_numScalarArguments + pCommandHeader->m_numOfLocalPtrSizes + ( pCommandHeader->m_numGlobalCapturedBuffer * ( sizeof( ulong ) +  sizeof( uint ) ) / sizeof( uint ) );

            pArgumentIds = &pCommandHeader->m_data[ offset ];
            //Object IDS are located after Argument IDs
            pObjectIds   = &pCommandHeader->m_data[ offset + pCommandHeader->m_numGlobalArguments ];

            //Setup local memory for fast access for Curbe patching
            uint ArgId = get_local_id( 0 );

            //Only third group Updates ObjectIDS, this will be synchronized with condition below
            if( ( ArgId >> HW_GROUP_ID_SHIFT( PARALLEL_SCHEDULER_COMPILATION_SIZE_20 ) ) == 2 )
            {
                ArgId = ArgId - ( PARALLEL_SCHEDULER_COMPILATION_SIZE_20 << 1 );
                while( ArgId < pCommandHeader->m_numGlobalArguments )
                {
                    objectIds[ pArgumentIds[ ArgId ] ] = pObjectIds[ ArgId ];
                    ArgId += PARALLEL_SCHEDULER_COMPILATION_SIZE_20;
                }
            }
#ifdef SCHEDULER_EMULATION
            //Synchronization needed for Emulation, ObjectIDS needs to be set by whole HW group, on GPU there is implicit synchronization in HW group
            barrier( CLK_GLOBAL_MEM_FENCE );
#endif
            if( get_local_id( 0 ) == PARALLEL_SCHEDULER_COMPILATION_SIZE_20 * 2 )
            {
                __global uint* pBindingTable = ( __global uint* ) &ssh[ btOffset ];


                //To properly set up binding table point to parents surface state heap
                for( uint ArgumentID = 0 ; ArgumentID < pCommandHeader->m_numGlobalArguments; ArgumentID++ )
                {
                    uint ArgId      = pArgumentIds[ ArgumentID ];

                    //Locate proper Arg ID
                    //Get ssh offset, lookup table already provided
                    if( objectIds[ ArgId ] < MAX_SSH_PER_KERNEL_SIZE )
                    {
                        if( pSSHdata[ ArgId ].m_sourceOffset == ArgId )
                        {
                            pBindingTable[ pSSHdata[ ArgId ].m_patchOffset ] = objectIds[ ArgId ];
                        }
                        else
                        {
                            pQueue->m_controls.m_ErrorCode += 10;
                            uint CurrentArg = 0;
                            while( CurrentArg < pKernelAddressData->m_BTSize / 4 )
                            {
                                if( pSSHdata[ CurrentArg ].m_sourceOffset == ArgId )
                                {
                                    pBindingTable[ pSSHdata[ CurrentArg ].m_patchOffset ] = objectIds[ ArgId ];
                                    break;
                                }
                                CurrentArg++;
                            }
                        }
                    }
                }
            }
        }

        if( ( PatchMask & SCHEDULER_DATA_PARAMETER_SAMPLER_MASK ) )
        {
            if( get_local_id( 0 ) == 2 * PARALLEL_SCHEDULER_COMPILATION_SIZE_20 )
            {
                for( uint ArgumentID = 0; ArgumentID < pCommandHeader->m_numGlobalArguments; ArgumentID++ )
                {
                    uint ArgId      = pArgumentIds[ ArgumentID ];
                    if( ( objectIds[ ArgId ] >= MAX_SSH_PER_KERNEL_SIZE ) )
                    {
                        uint SamplerCount = 0;
                        //Get pointer to Parent's samplers ( arguments ) data stored on KRS
                        __global IGIL_SamplerParams* pSamplerParamsOnKRS    = ( __global IGIL_SamplerParams* )GetPtrToKernelReflectionOffset( pKernelAddressData->m_SamplerParamsOffset, pKernelReflection );

                        //Iterate through all samplers passed from parent and copy state to proper SSA offset
                        while( pKernelReflection->m_ParentSamplerCount > SamplerCount )
                        {
                            //Get offset in parent's SSA from ObjectID, offset to beginning of SSA is included ( before SSA is BorderColorPointer ) so this is relative to parent's DSH heap
                            PatchOffset = objectIds[ ArgId ] - MAX_SSH_PER_KERNEL_SIZE;

                            if( pSamplerParamsOnKRS->m_ArgID == ArgId )
                            {
                                IGILLOCAL_MEMCPY_GTOG( &pDsh[ pSamplerParamsOnKRS->m_SamplerStateOffset ], &dsh[ pQueue->m_controls.m_ParentDSHOffset + PatchOffset ], OCLRT_SIZEOF_SAMPLER_STATE );
                                break;
                            }
                            pSamplerParamsOnKRS = pSamplerParamsOnKRS + 1;
                            SamplerCount        = SamplerCount + 1;
                        }
                    }
                }
            }
        }
        __global    char*   pScalarData                     = ( __global char* ) ( &pCommandHeader->m_data[ NumberOfDepencies ] );

        CurrentIndex = 0;
        uint TotalLocalSize = pWalkerEnum->WalkerArray[ WalkerID ].ActualLocalSize.x *
                              pWalkerEnum->WalkerArray[ WalkerID ].ActualLocalSize.y *
                              pWalkerEnum->WalkerArray[ WalkerID ].ActualLocalSize.z;

        uint NumberOfHWThreads = TotalLocalSize / pBlockData->m_SIMDSize;

        if( TotalLocalSize % pBlockData->m_SIMDSize != 0 )
        {
            NumberOfHWThreads++;
        }

        uint ThreadPayloadSize = NumberOfHWThreads * pBlockData->m_PayloadSize;

        //Move pointer to Constant Buffer Offset
        pDsh =  ( __global char* )&dsh[ DshOffset + SamplerHeapSize ];

        if( ( get_local_id( 0 ) >= PARALLEL_SCHEDULER_COMPILATION_SIZE_20 ) & ( get_local_id( 0 ) < PARALLEL_SCHEDULER_COMPILATION_SIZE_20 + 6 ) )
        {
            if( PatchMask & SCHEDULER_DATA_PARAMETER_KERNEL_ARGUMENT_MASK )
            {
                while( pKernelCurbeParams[ CurrentIndex ].m_parameterType == SCHEDULER_DATA_PARAMETER_KERNEL_ARGUMENT )
                {
                    CurbeSize               = pKernelCurbeParams[ CurrentIndex ].m_parameterSize;
                    PatchOffset             = pKernelCurbeParams[ CurrentIndex ].m_patchOffset;
                    IGILLOCAL_MEMCPY_GTOG( &pDsh[ PatchOffset ], pScalarData, CurbeSize );
                    pScalarData             += CurbeSize;
                    CurrentIndex++;
                }
            }
            if( PatchMask & SCHEDULER_DATA_PARAMETER_LOCAL_WORK_SIZE_MASK )
            {
                CurrentIndex     = PatchLocalWorkSizesParallel( CurrentIndex,
                                                                SCHEDULER_DATA_PARAMETER_LOCAL_WORK_SIZE,
                                                                pKernelCurbeParams,
                                                                pDsh,
                                                                LocalSizes.x,
                                                                LocalSizes.y,
                                                                LocalSizes.z,
                                                                pWalkerEnum->WalkerArray[ WalkerID ].ActualLocalSize.x,
                                                                pWalkerEnum->WalkerArray[ WalkerID ].ActualLocalSize.y,
                                                                pWalkerEnum->WalkerArray[ WalkerID ].ActualLocalSize.z );
            }
            if( PatchMask & SCHEDULER_DATA_PARAMETER_GLOBAL_WORK_SIZE_MASK )
            {
                CurrentIndex     = PatchDSH6TokensParallel20( CurrentIndex,
                                                              SCHEDULER_DATA_PARAMETER_GLOBAL_WORK_SIZE,
                                                              pKernelCurbeParams,
                                                              pDsh,
                                                              GlobalSizes.x,
                                                              GlobalSizes.y,
                                                              GlobalSizes.z );
            }
            if( PatchMask & SCHEDULER_DATA_PARAMETER_NUM_WORK_GROUPS_MASK )
            {
                CurrentIndex     = PatchDSH6TokensParallel20( CurrentIndex,
                                                              SCHEDULER_DATA_PARAMETER_NUM_WORK_GROUPS,
                                                              pKernelCurbeParams,
                                                              pDsh,
                                                              pWalkerEnum->TotalDimSize.x,
                                                              pWalkerEnum->TotalDimSize.y,
                                                              pWalkerEnum->TotalDimSize.z );
            }
            if( PatchMask & SCHEDULER_DATA_PARAMETER_WORK_DIMENSIONS_MASK )
            {
                CurrentIndex     = PatchDSH1TokenParallel20( CurrentIndex,
                                                             SCHEDULER_DATA_PARAMETER_WORK_DIMENSIONS,
                                                             pKernelCurbeParams,
                                                             pDsh,
                                                             WorkDim );
            }
            if( PatchMask & SCHEDULER_DATA_PARAMETER_SUM_OF_LOCAL_MEMORY_OBJECT_ARGUMENT_SIZES_MASK )
            {
                CurrentIndex     = PatchLocalMemEntities( CurrentIndex,
                                                          SCHEDULER_DATA_PARAMETER_SUM_OF_LOCAL_MEMORY_OBJECT_ARGUMENT_SIZES,
                                                          pKernelCurbeParams,
                                                          pDsh,
                                                          pCommandHeader );
            }
            if( PatchMask & SCHEDULER_DATA_PARAMETER_GLOBAL_WORK_OFFSET_MASK )
            {
                CurrentIndex     = PatchDSH6TokensParallel20( CurrentIndex,
                                                              SCHEDULER_DATA_PARAMETER_GLOBAL_WORK_OFFSET,
                                                              pKernelCurbeParams,
                                                              pDsh,
                                                              GlobalOffset.x,
                                                              GlobalOffset.y,
                                                              GlobalOffset.z );
            }
            if( PatchMask & SCHEDULER_DATA_PARAMETER_NUM_HARDWARE_THREADS_MASK )
            {
                CurrentIndex     = PatchDSH1TokenParallel20( CurrentIndex,
                                                             SCHEDULER_DATA_PARAMETER_NUM_HARDWARE_THREADS,
                                                             pKernelCurbeParams,
                                                             pDsh,
                                                             NumberOfHWThreads );
            }
            if( PatchMask & SCHEDULER_DATA_PARAMETER_PARENT_EVENT_MASK )
            {
                CurrentIndex     = PatchDSH1TokenParallel20( CurrentIndex,
                                                             SCHEDULER_DATA_PARAMETER_PARENT_EVENT,
                                                             pKernelCurbeParams,
                                                             pDsh,
                                                             pCommandHeader->m_event );
            }
            if( PatchMask & SCHEDULER_DATA_PARAMETER_ENQUEUED_LOCAL_WORK_SIZE_MASK )
            {
                CurrentIndex     = PatchDSH6TokensParallel20( CurrentIndex,
                                                              SCHEDULER_DATA_PARAMETER_ENQUEUED_LOCAL_WORK_SIZE,
                                                              pKernelCurbeParams,
                                                              pDsh,
                                                              LocalSizes.x,
                                                              LocalSizes.y,
                                                              LocalSizes.z );
            }
            if( PatchMask & SCHEDULER_DATA_PARAMETER_GLOBAL_POINTER )
            {
                if( pCommandHeader->m_numGlobalCapturedBuffer > 0 )
                {
                    //Handle global pointers patching in stateless mode, info about layout in declaration of IGIL_CommandHeader
                    __global    uint*  pGlobalIndexes = ( __global uint* ) ( &pCommandHeader->m_data[ NumberOfDepencies + pCommandHeader->m_numScalarArguments ] );
                    __global    uint*  pGlobalPtrs    = ( __global uint* ) ( &pCommandHeader->m_data[ NumberOfDepencies + pCommandHeader->m_numScalarArguments + pCommandHeader->m_numGlobalCapturedBuffer ] );
                    uint        StartIndex            = CurrentIndex;

                    //Argument in command header are not in correct sequence, that's why proper key needs to be located
                    for( uint glIdx = 0 ; glIdx < pCommandHeader->m_numGlobalCapturedBuffer; glIdx++)
                    {
                        //Reset CurrentIndex as we need to start from the beginning.
                        CurrentIndex  = StartIndex;
                        while( pKernelCurbeParams[ CurrentIndex ].m_parameterType == COMPILER_DATA_PARAMETER_GLOBAL_SURFACE )
                        {
                            //Patch only if exact match occurs
                            if( pKernelCurbeParams[ CurrentIndex ].m_sourceOffset == *pGlobalIndexes )
                            {
                                PatchOffset             = pKernelCurbeParams[ CurrentIndex ].m_patchOffset;
                                //64 bit patching
                                if( pKernelCurbeParams[ CurrentIndex ].m_parameterSize == 8 )
                                {
                                    __global uint* pDst = (__global uint *) &pDsh[PatchOffset];
                                    pDst[0] = pGlobalPtrs[0];
                                    pDst[1] = pGlobalPtrs[1];
                                }
                                else
                                {
                                    __global uint* pDst = ( __global uint* ) &pDsh[ PatchOffset ];
                                    *pDst               = ( uint ) *pGlobalPtrs;
                                }
                            }
                            CurrentIndex++;
                        }
                        pGlobalPtrs += 2;
                        pGlobalIndexes++;
                    }
                }
                while( pKernelCurbeParams[ CurrentIndex ].m_parameterType == COMPILER_DATA_PARAMETER_GLOBAL_SURFACE )
                {
                    CurrentIndex++;
                }
            }

            //Patch images curbe entries
            if( ( PatchMask & SCHEDULER_DATA_PARAMETER_IMAGE_CURBE_ENTRIES ) | ( PatchMask & SCHEDULER_DATA_PARAMETER_SAMPLER_MASK ) )
            {
                if( ( pArgumentIds != NULL ) & ( pObjectIds != NULL ) )
                {
                    //pKernelReflectionChar is a global address pointer
                    __global IGIL_ImageParamters      *pImageParams           = ( __global IGIL_ImageParamters * ) &pKernelReflectionChar[ pKernelReflection->m_ParentImageDataOffset ];
                    __global IGIL_ParentSamplerParams *pParentSamplerParams   = ( __global IGIL_ParentSamplerParams* ) &pKernelReflectionChar[ pKernelReflection->m_ParentSamplerParamsOffset ];
                    //First obtain argument ID
                    uint WorkID = get_local_id( 0 ) - PARALLEL_SCHEDULER_COMPILATION_SIZE_20;
                    while( WorkID + CurrentIndex < pBlockData->m_numberOfCurbeTokens )
                    {
                        uint ArgId              = pKernelCurbeParams[ CurrentIndex + WorkID ].m_sourceOffset;
                        uint ObjectID           = objectIds[ ArgId ];
                        uint CurrentImage       = 0;
                        uint CurrentSampler     = 0;

                        uint PatchValue         = 0;
                        uint TokenType          = pKernelCurbeParams[ CurrentIndex + WorkID ].m_parameterType;
                        uint PatchOffset        = pKernelCurbeParams[ CurrentIndex + WorkID ].m_patchOffset;
                        uint PatchValueInvalid  = 0;

                        //If Images
                        if( ObjectID < OCLRT_IMAGE_MAX_OBJECT_ID )
                        {
                            //Locate proper parent Image
                            while( ( pImageParams[ CurrentImage ].m_ObjectID != ObjectID ) & ( CurrentImage < pKernelReflection->m_ParentKernelImageCount ) )
                            {
                                CurrentImage++;
                            }
                            //Proper image is located under CurrentImage patch the token

                            if( TokenType == SCHEDULER_DATA_PARAMETER_IMAGE_WIDTH )
                            {
                                PatchValue  = pImageParams[ CurrentImage ].m_Width;
                            }
                            else if( TokenType == SCHEDULER_DATA_PARAMETER_IMAGE_HEIGHT )
                            {
                                PatchValue  = pImageParams[ CurrentImage ].m_Height;
                            }
                            else if( TokenType == SCHEDULER_DATA_PARAMETER_IMAGE_DEPTH )
                            {
                                PatchValue  = pImageParams[ CurrentImage ].m_Depth;
                            }
                            else if( TokenType == SCHEDULER_DATA_PARAMETER_IMAGE_CHANNEL_DATA_TYPE )
                            {
                                PatchValue  = pImageParams[ CurrentImage ].m_ChannelDataType;
                            }
                            else if( TokenType == SCHEDULER_DATA_PARAMETER_IMAGE_CHANNEL_ORDER )
                            {
                                PatchValue  = pImageParams[ CurrentImage ].m_ChannelOrder;
                            }
                            else if( TokenType == SCHEDULER_DATA_PARAMETER_IMAGE_SRGB_CHANNEL_ORDER )
                            {
                                PatchValue  = GetSRGBChannelOrder( pImageParams[ CurrentImage ].m_ChannelOrder );
                            }
                            else if( TokenType == SCHEDULER_DATA_PARAMETER_IMAGE_ARRAY_SIZE )
                            {
                                PatchValue  = pImageParams[ CurrentImage ].m_ArraySize;
                            }
                            else if( TokenType == SCHEDULER_DATA_PARAMETER_IMAGE_NUM_SAMPLES )
                            {
                                PatchValue  = pImageParams[ CurrentImage ].m_NumSamples;
                            }
                            else if( TokenType == SCHEDULER_DATA_PARAMETER_IMAGE_NUM_MIP_LEVELS )
                            {
                                PatchValue  = pImageParams[ CurrentImage ].m_NumMipLevels;
                            }
                            else if( TokenType == SCHEDULER_DATA_PARAMETER_IMAGE_OBJECT_ID )
                            {
                                PatchValue  = ObjectID;
                            }
                            else
                            {
                                PatchValueInvalid = 1;
                            }
                        }
                        //If Sampler
                        else if( ObjectID >= OCLRT_SAMPLER_MIN_OBJECT_ID )
                        {
                            //Mark PatchValue invalid if SamplerParams will not be found
                            PatchValueInvalid = 1;
                            //Locate proper parent Image
                            while( CurrentSampler < pKernelReflection->m_ParentSamplerCount )
                            {
                                if( pParentSamplerParams[ CurrentSampler ].m_ObjectID == ObjectID )
                                {
                                    PatchValueInvalid = 0;
                                    if( TokenType == DATA_PARAMETER_SAMPLER_ADDRESS_MODE )
                                    {
                                        PatchValue = pParentSamplerParams[ CurrentSampler ].m_AddressingMode;
                                    }
                                    else if( TokenType == DATA_PARAMETER_SAMPLER_NORMALIZED_COORDS )
                                    {
                                        PatchValue = pParentSamplerParams[ CurrentSampler ].NormalizedCoords;
                                    }
                                    else if( TokenType == DATA_PARAMETER_SAMPLER_COORDINATE_SNAP_WA_REQUIRED )
                                    {
                                        PatchValue = pParentSamplerParams[ CurrentSampler ].CoordinateSnapRequired;
                                    }
                                    else if( TokenType == SCHEDULER_DATA_PARAMETER_SAMPLER_OBJECT_ID )
                                    {
                                        PatchValue = ObjectID;
                                    }
                                    else
                                    {
                                        PatchValueInvalid = 1;
                                    }
                                    CurrentSampler = pKernelReflection->m_ParentSamplerCount;
                                }
                                CurrentSampler++;
                            }
                        }
                        else
                        {
                            PatchValueInvalid = 1;
                        }

                        if( PatchValueInvalid == 0 )
                        {
                            *( __global uint * ) ( &pDsh[ PatchOffset ] ) = PatchValue;
                        }
                        CurrentIndex += 6;
                    }

                }
                else
                {
                    pQueue->m_controls.m_ErrorCode += 7;
                }
            }
        }
#ifdef SCHEDULER_EMULATION
        barrier( CLK_GLOBAL_MEM_FENCE );
#endif

        if( get_local_id( 0 ) == 0 )
        {
#if defined WA_LRI_COMMANDS_EXIST
            bool ShouldDisablePreemption = false;
#endif
            //Profiling support
            if( pQueue->m_controls.m_IsProfilingEnabled != 0 )
            {
                bool DisableTimeStampStart = true;
                bool DisableTimeStampEnd   = true;
                if( ( ( uint )pCommandHeader->m_event != IGIL_EVENT_INVALID_HANDLE ) & ( ( WalkerID == 0 ) | ( WalkerID == WalkerCount - 1 ) ) )
                {
                    //Event is propagated to childs as "parent event", to avoid overwriting the same start value, only generate timestamp write
                    //For the first command for this event, this means we look for event with no children ( so compare to 1 ).
                    clk_event_t EventID = __builtin_astype( ( void* ) ( ( ulong ) pCommandHeader->m_event ), clk_event_t );
                    __global IGIL_DeviceEvent *events = TEMP_IGIL_GetDeviceEvents( eventsPool );

                    if( events[ ( uint )(size_t)__builtin_astype( EventID, void* ) ].m_numChildren == 1 )
                    {
#if defined WA_LRI_COMMANDS_EXIST && defined WA_PROFILING_PREEMPTION
                        //This is a case, where profiling of block kernels occurs - presence of event in EM workload
                        //In such case, disable preemption around all WALKERs for that block kernel and event
                        ShouldDisablePreemption = true;
#endif
                        if( WalkerID == 0 )
                        {
                            //Emit pipecontrol with timestamp write
                            ulong Address = ( ulong )&( events[ ( uint )(size_t)__builtin_astype( EventID, void* ) ].m_profilingCmdStart );
                            //Timestamp start
                            PatchPipeControlProfilingAddres( secondaryBatchBuffer,
                                                             SLBOffset,
                                                             Address,
                                                             PIPE_CONTROL_FOR_TIMESTAMP_START_OFFSET );
                            DisableTimeStampStart = false;
                        }
                        if( WalkerID == WalkerCount - 1 )
                        {
                            ulong Address = ( ulong )&( events[ ( uint )(size_t)__builtin_astype( EventID, void* ) ].m_profilingCmdEnd );
                            //Timestamp end
                            PatchPipeControlProfilingAddres( secondaryBatchBuffer,
                                                             SLBOffset,
                                                             Address,
                                                             PIPE_CONTROL_FOR_TIMESTAMP_END_OFFSET );
                            DisableTimeStampEnd = false;
                        }
                    }
                }
                if( DisableTimeStampStart )
                {
                    DisablePostSyncBitInPipeControl( secondaryBatchBuffer,
                                                     SLBOffset,
                                                     PIPE_CONTROL_FOR_TIMESTAMP_START_OFFSET );
                }
                if( DisableTimeStampEnd )
                {
                    DisablePostSyncBitInPipeControl( secondaryBatchBuffer,
                                                     SLBOffset,
                                                     PIPE_CONTROL_FOR_TIMESTAMP_END_OFFSET );
                }
            }
            else
            {
                //Optimized path, in case block can be run concurently noop pipe control after such block.
                uint DwordOffset = SLBOffset / DWORD_SIZE_IN_BYTES;

                if( pBlockData->m_CanRunConcurently != 0 )
                {
                    NOOPCSStallPipeControl( secondaryBatchBuffer, DwordOffset, PIPE_CONTROL_FOR_TIMESTAMP_END_OFFSET );
                }
                else
                {
                    PutCSStallPipeControl( secondaryBatchBuffer, DwordOffset, PIPE_CONTROL_FOR_TIMESTAMP_END_OFFSET );
                }
            }

#if defined WA_LRI_COMMANDS_EXIST
            bool NoopPreemptionDisabling = true;
            bool NoopPreemptionEnabling = true;

#if defined WA_KERNEL_PREEMPTION
            //This is case, where block kernel should have disabled preemption because of its sampler usage around all WALKERs of that block kernel
            //Preemption should be disabled when EM event profiling is used OR kernel data indicate such behavior
            ShouldDisablePreemption |= ( pBlockData->m_DisablePreemption != 0 );
#endif

#if defined WA_PROFILING_PREEMPTION
            //m_EventTimestampAddress != NULL means profiling of the whole workload is enabled (preemption around whole chained BB is disabled)
            //So disabling preemption should be permitted only when workload profiling is off, in other cases noop all LRI commands
            //For m_EventTimestampAddress != NULL preemption is enabled before BB_END
            ShouldDisablePreemption &= ( pQueue->m_controls.m_EventTimestampAddress == 0 );
#endif

            if( ShouldDisablePreemption != false )
            {
                if( WalkerID == 0 )
                {
                    SetDisablePreemptionRegister( SLBOffset, secondaryBatchBuffer );
                    NoopPreemptionDisabling = false;
                }

                if( WalkerID == WalkerCount - 1 )
                {
                    SetEnablePreemptionRegister( SLBOffset, secondaryBatchBuffer );
                    NoopPreemptionEnabling = false;
                }
            }

            if( NoopPreemptionDisabling )
            {
                NoopPreemptionCommand( SLBOffset, IMM_LOAD_REGISTER_FOR_DISABLE_PREEMPTION_OFFSET, secondaryBatchBuffer );
            }

            if( NoopPreemptionEnabling )
            {
                NoopPreemptionCommand( SLBOffset, IMM_LOAD_REGISTER_FOR_ENABLE_PREEMPTION_OFFSET, secondaryBatchBuffer );
            }
#endif //WA_LRI_COMMANDS_EXIST
        }

        //Witems from 0 to 16 are responsible for local ids generation.
        if( ( get_local_id( 0 ) < 16 ) & ( pBlockData->m_NeedLocalIDS != 0 ) )
        {
            //Now generate local IDS
            generateLocalIDSParallel20( pLocalIdsOnDSH, pWalkerEnum->WalkerArray[ WalkerID ].ActualLocalSize, NumberOfHWThreads, pBlockData->m_SIMDSize );
        }
        //3rd HW thread will take care of patching media curbe load and GPPGU_WALKER command
        if( get_local_id( 0 ) == PARALLEL_SCHEDULER_COMPILATION_SIZE_20 * 2 )
        {

            uint TotalSLMSize = pCommandHeader->m_totalLocalSize + pBlockData->m_InilineSLMSize;
            //Update Interface Descriptor Data with SLM size  / number of HW threads.
            CopyAndPatchIDData20(dsh, blockId, NumberOfHWThreads, TotalSLMSize, IntefaceDescriptorOffset,
                                 pQueue->m_controls.m_StartBlockID, btOffset, DshOffset,
                                 pBlockData->m_numberOfSamplerStates
#ifdef ENABLE_DEBUG_BUFFER
                                 ,
                                 DebugQueue
#endif
            );

            patchGpGpuWalker( SLBOffset, secondaryBatchBuffer, IntefaceDescriptorOffset, pBlockData->m_SIMDSize,
                              TotalLocalSize, pWalkerEnum->WalkerArray[ WalkerID ].WalkerDimSize, pWalkerEnum->WalkerArray[ WalkerID ].WalkerStartPoint,
                              NumberOfHWThreads, pBlockData->m_sizeOfConstantBuffer + ThreadPayloadSize, SamplerHeapSize + DshOffset );

            PatchMediaStateFlush( SLBOffset, secondaryBatchBuffer, IntefaceDescriptorOffset, SCHEDULER_MSF_INITIAL );
            PatchMediaStateFlush( SLBOffset, secondaryBatchBuffer, IntefaceDescriptorOffset, SCHEDULER_MSF_SECOND );
        }
    }
}

uint CheckEventStatus( __global IGIL_CommandHeader* pCommand,
                       __global IGIL_EventPool* eventsPool )
{
    if( pCommand->m_numDependencies == 0 )
    {
        return 0;
    }
    else
    {
         __global IGIL_DeviceEvent* pDeviceEvent;
         //Events are stored at the begining of command packet dynamic payload
         for( uint i = 0; i < pCommand->m_numDependencies; i++ )
         {
             pDeviceEvent = TEMP_IGIL_GetDeviceEvent( eventsPool, pCommand->m_data[ i ] );
             if( pDeviceEvent->m_state != CL_COMPLETE )
             {
                 return 1;
             }
         }
    }
    return 0;
}

void DecreaseEventDependenciesParallel( __global IGIL_CommandHeader* pCommand,
                                        __global IGIL_EventPool* eventsPool )
{
     __global IGIL_DeviceEvent* pDeviceEvent;

     //Events are stored at the begining of command packet dynamic payload
     for( uint i = 0; i < pCommand->m_numDependencies; i++ )
     {
         pDeviceEvent = TEMP_IGIL_GetDeviceEvent( eventsPool, pCommand->m_data[ i ] );
         int OldDependants = atomic_dec( &pDeviceEvent->m_numDependents );

         if( ( pDeviceEvent->m_refCount <= 0 ) &
             ( ( OldDependants - 1 ) <= 0 ) &
             ( pDeviceEvent->m_numChildren <= 0 ) )
             {
                 TEMP_IGIL_FreeEvent( __builtin_astype( ( void* )( ( ulong )pCommand->m_data[ i ] ), clk_event_t ), eventsPool );
             }
    }
}

//Update status of the event and all events that are depending on this event
void UpdateEventsTreeStatusParallel( clk_event_t eventId, __global IGIL_EventPool* eventsPool, bool isProfilingEnabled )
{
    __global IGIL_DeviceEvent *events = TEMP_IGIL_GetDeviceEvents( eventsPool );
    __global IGIL_DeviceEvent *pEvent;
    do
    {
        pEvent = &events[ (uint) (size_t)__builtin_astype( eventId, void* ) ];

        int OldNumChild = atomic_dec( &pEvent->m_numChildren );
        if( ( OldNumChild - 1 ) <= 0 )
        {
            pEvent->m_state = CL_COMPLETE;

            if( ( pEvent->m_refCount <= 0 ) &
                ( pEvent->m_numDependents <= 0 ) &
                ( pEvent->m_numChildren <= 0 ) )
            {
                TEMP_IGIL_FreeEvent( eventId, eventsPool );
            }
            //This event transitions to CL_COMPLETE state, update it profiling informations.
            if( isProfilingEnabled != 0 )
            {
                //CL COMPLETE time is before this scheduler starts
                pEvent->m_profilingCmdComplete = eventsPool->m_CLcompleteTimestamp;

                //Check if this event has profiling pointer, if so update profiling data, all times should be there atm
                if( pEvent->m_pProfiling != 0 )
                {
                    __global ulong* retValues = ( __global ulong * )pEvent->m_pProfiling;

                    ulong StartTime                = pEvent->m_profilingCmdStart;
                    ulong EndTime                  = pEvent->m_profilingCmdEnd;
                    ulong CompleteTime             = pEvent->m_profilingCmdComplete;
                    ulong CLEndTransitionTime      = 0;
                    ulong CLCompleteTransitionTime = 0;
                    //Check if timer didn't reset by hitting max value
                    if( CompleteTime > StartTime )
                    {
                        CLEndTransitionTime      = EndTime - StartTime;
                        CLCompleteTransitionTime = CompleteTime - StartTime;
                    }
                    //If we hit this else it means that GPU timer reset to 0, compute proper delta
                    else
                    {
                        if( EndTime < StartTime )
                        {
                            CLEndTransitionTime = PROFILING_MAX_TIMER_VALUE - StartTime + EndTime;
                        }
                        else
                        {
                            CLEndTransitionTime = EndTime - StartTime;
                        }
                        CLCompleteTransitionTime = PROFILING_MAX_TIMER_VALUE - StartTime + CompleteTime;
                    }
                    //First value is END - START timestamp
                    retValues[ 0 ] = ( ulong )( ( float )CLEndTransitionTime * eventsPool->m_TimestampResolution );
                    //Second value is COMPLETE - START timestamp
                    retValues[ 1 ] = ( ulong )( ( float )CLCompleteTransitionTime * eventsPool->m_TimestampResolution );
                }
            }
            //Signal parent because we completed
            eventId = __builtin_astype( ( void* )( ( ulong )pEvent->m_parentEvent ), clk_event_t );
        }
    }
    while ( ( ( uint )(size_t)__builtin_astype( eventId, void* ) != IGIL_EVENT_INVALID_HANDLE ) & ( pEvent->m_numChildren <= 0 ) );
}

void GlobalBarrier( __global volatile uint* syncSurface )
{
    //Make sure each WKG item hit the barrier.
    barrier( CLK_GLOBAL_MEM_FENCE );

    //Now first thread of each wkg writes to designated place on SyncSurface
    if ( get_local_id( 0 ) == 0 )
    {
        syncSurface[ get_group_id( 0 ) ] = 1;
    }
    //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( get_group_id( 0 ) == ( get_num_groups( 0 ) - 1 ) )
    {
        //24 -48 case
        uint Value;
        do
        {
            Value = 1;
            for( uint i = get_local_id( 0 ); i < get_num_groups( 0 ); i += get_local_size( 0 ) )
            {
                Value = Value & syncSurface[ i ];
            }

        }
        while( Value == 0 );
        barrier( CLK_GLOBAL_MEM_FENCE );

        for( uint i = get_local_id( 0 ); i < get_num_groups( 0 ); i += get_local_size( 0 ) )
        {
            syncSurface[ i ] = 0;
        }
    }

    if( get_local_id( 0 ) == 0 )
    {
        while( syncSurface[ get_group_id( 0 ) ] != 0 );
    }
    barrier( CLK_GLOBAL_MEM_FENCE );
}

void GlobalBarrierUpdateQueue( __global volatile uint* syncSurface, __global IGIL_CommandQueue* pQueue )
{
    //Make sure each WKG item hit the barrier.
    barrier( CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE);
    //Now first thread of each wkg writes to designated place on SyncSurface
    if ( get_local_id(0) == 0 )
    {
        syncSurface[get_group_id(0)] = 1;
    }
    //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( get_group_id(0) == ( get_num_groups( 0 ) - 1 ) )
    {
        uint Value;
        do
        {
            Value = 1;
            for( uint i = get_local_id( 0 ); i < get_num_groups( 0 ); i += get_local_size( 0 ) )
            {
                Value = Value & syncSurface[ i ];
            }
        }
        while( Value == 0 );
        barrier( CLK_GLOBAL_MEM_FENCE );
        pQueue->m_controls.m_IDTAfterFirstPhase = pQueue->m_controls.m_CurrentIDToffset;

        barrier( CLK_GLOBAL_MEM_FENCE );

        for( uint i = get_local_id( 0 ); i < get_num_groups( 0 ); i += get_local_size( 0 ) )
        {
            syncSurface[ i ] = 0;
        }
    }

    if( get_local_id(0) == 0 )
    {
        while( syncSurface[ get_group_id(0) ] != 0 );
    }
    barrier( CLK_GLOBAL_MEM_FENCE );
}


#ifdef SCHEDULER_EMULATION
__local int IDTOffset;
__local int DSHOffset;
__local int SLBOffset;
__local int StackOffset;
__local int QStorageOffset;
__local int MarkerOffset;
__local int BTSoffset;
__local IGIL_WalkerEnumeration WalkerEnum;
__local uint ObjectIDS[ MAX_GLOBAL_ARGS ];
#endif

#define WA_INT_DESC_MAX 62
__kernel __attribute__((intel_reqd_sub_group_size(PARALLEL_SCHEDULER_COMPILATION_SIZE_20)))
void SchedulerParallel20(
    __global IGIL_CommandQueue* pQueue,
    __global uint* commandsStack,
    __global IGIL_EventPool*    eventsPool,
    __global uint* secondaryBatchBuffer,        //SLB that will be used to put commands in.
    __global char* dsh,                         //Pointer to the start of Dynamic State Heap
    __global IGIL_KernelDataHeader* kernelData, //This is kernel reflection surface
    __global volatile uint* queueStorageBuffer,
    __global char* ssh,                         //Pointer to Surface state heap with BT and SS
    __global DebugDataBuffer* debugQueue )
{
    EMULATION_ENTER_FUNCTION( );
#ifdef WA_DISABLE_SCHEDULERS
    return;
#endif

#ifdef DEBUG
    //Early return enabled when m_SchedulerEarlyReturn is > 0,
    if( pQueue->m_controls.m_SchedulerEarlyReturn > 0 )
    {
        if( pQueue->m_controls.m_SchedulerEarlyReturn == 1 )
        {
            return;
        }

        if( get_global_id( 0 ) == 0 )
        {
            pQueue->m_controls.m_SchedulerEarlyReturnCounter++;
        }
        GlobalBarrier( queueStorageBuffer );

        if( pQueue->m_controls.m_SchedulerEarlyReturnCounter == pQueue->m_controls.m_SchedulerEarlyReturn )
        {
            if( ( ( get_group_id( 0 ) == 1) == ( get_num_groups( 0 ) > 1 )  ) & ( get_local_id( 0 ) == 0 ) )
            {
#ifdef ENABLE_DEBUG_BUFFER
                //Set START time of current (last) scheduler
                if( ( pQueue->m_controls.m_IsProfilingEnabled != 0 ) & ( DebugQueue != 0 ) & ( DebugQueue->m_flags == DDB_SCHEDULER_PROFILING ) )
                {
                    *( ( __global ulong * ) ( &DebugQueue->m_data[ atomic_add( &DebugQueue->m_offset, 2 ) ] ) ) = EventsPool->m_CLcompleteTimestamp;
                }
#endif
                pQueue->m_controls.Temporary[ 2 ]++;

                //SlbOffset is expressed in bytes and for cmd it is needed to convert it to dwords
                __private uint DwordOffset = ( pQueue->m_controls.m_SecondLevelBatchOffset % MAX_SLB_OFFSET ) / DWORD_SIZE_IN_BYTES;
                //BB_START 1st DWORD
                secondaryBatchBuffer[ DwordOffset ] = OCLRT_BATCH_BUFFER_BEGIN_CMD_DWORD0;
                DwordOffset++;
                //BB_START 2nd DWORD - Address, 3rd DWORD Address high
                secondaryBatchBuffer[ DwordOffset++ ] = (uint)(pQueue->m_controls.m_CleanupSectionAddress & 0xFFFFFFFF);
                secondaryBatchBuffer[ DwordOffset ] = (uint)((pQueue->m_controls.m_CleanupSectionAddress >> 32) & 0xFFFFFFFF);
            }
            return;
        }
    }
#endif

    //First check if there are any new command packets on queue_t
    __global IGIL_CommandHeader* pCommand = 0;
    uint GroupID                          = get_group_id( 0 );

#ifndef SCHEDULER_EMULATION
    __local int IDTOffset;
    __local int DSHOffset;
    __local int SLBOffset;
    __local int StackOffset;
    __local int QStorageOffset;
    __local int MarkerOffset;
    __local int BTSoffset;
    __local IGIL_WalkerEnumeration WalkerEnum;
    __local uint ObjectIDS[ MAX_GLOBAL_ARGS ];

#endif

    if( pQueue->m_controls.m_LastScheduleEventNumber > 0 )
    {
        //Check if there are any events that needs updating, each wkg uses all hw threads in wkg to update events
        if( GroupID * PARALLEL_SCHEDULER_HWTHREADS_IN_HW_GROUP20 < pQueue->m_controls.m_LastScheduleEventNumber )
        {
            clk_event_t EventID;
            if( get_local_id( 0 ) % PARALLEL_SCHEDULER_COMPILATION_SIZE_20 == 0 )
            {
                uint ID = ( GroupID * PARALLEL_SCHEDULER_HWTHREADS_IN_HW_GROUP20 ) + ( get_local_id( 0 ) / PARALLEL_SCHEDULER_COMPILATION_SIZE_20 );
                while( ID < pQueue->m_controls.m_LastScheduleEventNumber )
                {
                    EventID = __builtin_astype( ( void* )( ( ulong )pQueue->m_controls.m_EventDependencies[ ID ] ), clk_event_t );
                    UpdateEventsTreeStatusParallel( EventID, eventsPool, ( pQueue->m_controls.m_IsProfilingEnabled != 0 ) );
                    ID      += get_num_groups( 0 ) * PARALLEL_SCHEDULER_HWTHREADS_IN_HW_GROUP20;
                }
            }
        }
        GlobalBarrier( queueStorageBuffer );
    }
    //Queue parsing section
    uint NumberOfEnqueues = pQueue->m_controls.m_TotalNumberOfQueues - pQueue->m_controls.m_PreviousNumberOfQueues;
    if( NumberOfEnqueues > 0 )
    {
        uint InitialOffset              = pQueue->m_controls.m_PreviousHead;
        bool PacketScheduled            = true;
        uint offset                     = 0;

        for( uint CurrentPacket = GroupID; CurrentPacket < NumberOfEnqueues; CurrentPacket += get_num_groups( 0 ) )
        {
            if( CurrentPacket == GroupID )
            {
                offset                  = TEMP_IGIL_GetNthCommandHeaderOffset( pQueue, InitialOffset, CurrentPacket );
            }
            else
            {
                offset                  = TEMP_IGIL_GetNthCommandHeaderOffset( pQueue, offset, get_num_groups( 0 ) );
            }
            pCommand                    = TEMP_IGIL_GetCommandHeader( pQueue, offset );

            //Initialize command packet with proper lws
            if( get_local_id( 0 ) == 0 )
            {
                InitializeCommandPacket( pCommand );
            }

            //Can I run this command ?
            if( CheckEventStatus( pCommand, eventsPool ) == 0 )
            {
                //Is it marker command ?
                if( pCommand->m_kernelId != IGIL_KERNEL_ID_ENQUEUE_MARKER )
                {
                    //Is there enough IDT space for me ?
                    if( get_local_id( 0 ) == 0 )
                    {
                        int WalkerNeeded = GetWalkerCount( pCommand );
                        //Optimization - check if IDT has free space for me
                        if( pQueue->m_controls.m_CurrentIDToffset + WalkerNeeded  <= WA_INT_DESC_MAX )
                        {
                            uint Temp = atomic_add( &pQueue->m_controls.m_CurrentIDToffset, WalkerNeeded );
                            if( Temp + WalkerNeeded <= WA_INT_DESC_MAX )
                            {
                                IDTOffset = Temp;
                                DSHOffset = atomic_add( &pQueue->m_controls.m_CurrentDSHoffset, ( MAX_DSH_SIZE_PER_ENQUEUE * WalkerNeeded ) );
                                SLBOffset = ( ( atomic_add( &pQueue->m_controls.m_SecondLevelBatchOffset, ( SECOND_LEVEL_BUFFER_SPACE_FOR_EACH_ENQUEUE * WalkerNeeded ) ) ) % MAX_SLB_OFFSET );
                                BTSoffset = atomic_add( &pQueue->m_controls.m_CurrentSSHoffset, pQueue->m_controls.m_BTmaxSize );
                            }
                            else
                            {
                                IDTOffset = -1;
                            }
                        }
                        else
                        {
                            IDTOffset = -1;
                        }
                    }
                    //Now barrier and check if we can go with scheduling
                    barrier( CLK_LOCAL_MEM_FENCE );

                    if( IDTOffset != -1 )
                    {
                        //This packet is all set, schedule it and we are done with it.
                        //Patch DSH has media curbe load and patch gpgpu walker inside
                        PatchDSHParallelWithDynamicDSH20( SLBOffset,
                                                          DSHOffset,
                                                          IDTOffset,
                                                          kernelData,
                                                          dsh,
                                                          pCommand->m_kernelId,
                                                          pCommand,
                                                          secondaryBatchBuffer,
                                                          pQueue,
                                                          eventsPool,
                                                          ssh,
                                                          BTSoffset,
                                                          &WalkerEnum,
                                                          ObjectIDS
#ifdef ENABLE_DEBUG_BUFFER
                                                          , DebugQueue
#endif
                                                          );
                        PacketScheduled = true;
                    }
                    else
                    {
                        PacketScheduled = false;
                    }
                }
                else //For marker we need to update returned event status
                {
                    //Check if there is space to track event
                    if( get_local_id( 0 ) == 0 )
                    {
                        uint Temp = atomic_inc( &pQueue->m_controls.m_EnqueueMarkerScheduled );
                        if( Temp < MAX_NUMBER_OF_ENQUEUE_MARKER )
                        {
                            MarkerOffset = Temp;
                        }
                        else
                        {
                            MarkerOffset = -1;
                        }
                    }
                    barrier( CLK_LOCAL_MEM_FENCE );
                    if( MarkerOffset != -1 )
                    {
                        PacketScheduled = true;
                    }
                    else
                    {
                        PacketScheduled = false;
                    }
                }
                //Update event dependencies if any, if there are event waiting for status change, put them on the list.
                if( PacketScheduled == true )
                {
                    if( get_local_id( 0 ) == 0 )
                    {
                        if( ( uint )pCommand->m_event != IGIL_EVENT_INVALID_HANDLE )
                        {
                           pQueue->m_controls.m_EventDependencies[ atomic_inc( &pQueue->m_controls.m_CurrentScheduleEventNumber ) ] = pCommand->m_event;
                        }
                        //Remove event dependencies setting.
                        if( pCommand->m_numDependencies > 0 )
                        {
                            DecreaseEventDependenciesParallel( pCommand, eventsPool );
                        }
                    }
                }
            }
            //Can't schedule it right now, move to storage.
            else
            {
                if( pQueue->m_controls.m_IsSimulation )
                {
                    barrier( CLK_LOCAL_MEM_FENCE );
                }
                PacketScheduled = false;
            }

            //Allocation failure, move command to stack storage and update stack pointers
            if( PacketScheduled == false )
            {
                if( get_local_id( 0 ) == 0 )
                {
                    StackOffset                  = atomic_dec( &pQueue->m_controls.m_StackTop ) - 1;
                    QStorageOffset               = atomic_sub( &pQueue->m_controls.m_QstorageTop, pCommand->m_commandSize ) - pCommand->m_commandSize;
                    commandsStack[ StackOffset ] = QStorageOffset;
                }
                barrier( CLK_LOCAL_MEM_FENCE );
                __global char* ptrQueue = ( __global char* )pQueue;
                GLOBAL_MEMCPY( ( __global void* )&queueStorageBuffer[ QStorageOffset / 4 ], (__global void * )&ptrQueue[ offset ] , pCommand->m_commandSize );
            }
            else if( pQueue->m_controls.m_IsSimulation )
            {
                barrier( CLK_LOCAL_MEM_FENCE );
            }
        }

        //In case there were new enqueues on queue_t, all work items must hit the global barrier before they can start taking items from the stack.
        GlobalBarrierUpdateQueue( queueStorageBuffer, pQueue );
    }

    //Check stack only when there are free IDTS
    if( ( pQueue->m_controls.m_IDTAfterFirstPhase < WA_INT_DESC_MAX ) &
        ( pQueue->m_controls.m_PreviousStackTop != pQueue->m_controls.m_StackSize ) )
    {
        //Start stack browsing
        uint MyID            = get_group_id( 0 );
        //Start browsing from the begining of the previous stack top
        uint CurrentOffset   = pQueue->m_controls.m_PreviousStackTop + MyID;
        uint CommandOffset   = 0;

        while( CurrentOffset < pQueue->m_controls.m_StackSize )
        {
            CommandOffset = commandsStack[ CurrentOffset ];

            if( CommandOffset != 0 )
            {
                pCommand = GetCommandHeaderFromStorage( ( __global uint* )queueStorageBuffer, CommandOffset );

                //Can I run this command ?
                if( CheckEventStatus( pCommand, eventsPool ) == 0 )
                {
                    //Is it marker command ?
                    if( pCommand->m_kernelId != IGIL_KERNEL_ID_ENQUEUE_MARKER )
                    {
                        //Is there enough IDT space for me ?
                        if( get_local_id( 0 ) == 0 )
                        {
                            int WalkerNeeded = GetWalkerCount( pCommand );
                            //Optimization - check if IDT has free space for me
                            if( pQueue->m_controls.m_CurrentIDToffset + WalkerNeeded  <= WA_INT_DESC_MAX )
                            {
                                uint Temp = atomic_add( &pQueue->m_controls.m_CurrentIDToffset, WalkerNeeded );
                                if( Temp + WalkerNeeded <= WA_INT_DESC_MAX )
                                {
                                    IDTOffset = Temp;
                                    DSHOffset = atomic_add( &pQueue->m_controls.m_CurrentDSHoffset, ( MAX_DSH_SIZE_PER_ENQUEUE * WalkerNeeded ) );
                                    SLBOffset = ( ( atomic_add( &pQueue->m_controls.m_SecondLevelBatchOffset, ( SECOND_LEVEL_BUFFER_SPACE_FOR_EACH_ENQUEUE * WalkerNeeded ) ) ) % MAX_SLB_OFFSET );
                                    BTSoffset = atomic_add( &pQueue->m_controls.m_CurrentSSHoffset, pQueue->m_controls.m_BTmaxSize );
                                }
                                else
                                {
                                    IDTOffset = -1;
                                }
                            }
                            else
                            {
                                 IDTOffset = -1;
                            }
                        }
                        //Now barrier and check if we can go with scheduling
                        barrier( CLK_LOCAL_MEM_FENCE );
                        if( IDTOffset != -1 )
                        {
                            //This packet is all set, schedule it and we are done with it.
                            //Patch DSH has media curbe load and patch gpgpu walker inside
                            PatchDSHParallelWithDynamicDSH20( SLBOffset,
                                                              DSHOffset,
                                                              IDTOffset,
                                                              kernelData,
                                                              dsh,
                                                              pCommand->m_kernelId,
                                                              pCommand,
                                                              secondaryBatchBuffer,
                                                              pQueue,
                                                              eventsPool,
                                                              ssh,
                                                              BTSoffset,
                                                              &WalkerEnum,
                                                              ObjectIDS
#ifdef ENABLE_DEBUG_BUFFER
                                                              , DebugQueue
#endif
                                                              );
                            pCommand->m_commandState = CAN_BE_RECLAIMED;
                            //Reset stack offset
                            commandsStack[ CurrentOffset ] = 0;

                            //Update event status
                            if( get_local_id( 0 ) == 0 )
                            {
                                //Add events dependant on this command to list of events neeeded to be updated.
                                if( ( uint )pCommand->m_event != IGIL_EVENT_INVALID_HANDLE )
                                {
                                    pQueue->m_controls.m_EventDependencies[ atomic_inc( &pQueue->m_controls.m_CurrentScheduleEventNumber ) ] = pCommand->m_event;
                                }

                                //Remove event dependencies setting.
                                if( pCommand->m_numDependencies > 0 )
                                {
                                    DecreaseEventDependenciesParallel( pCommand, eventsPool );
                                }
                            }
                        }
                    }
                    else // For marker we need to update returned event status
                    {
                        barrier( CLK_GLOBAL_MEM_FENCE );
                        //Check if there is space to track event
                        if( get_local_id( 0 ) == 0 )
                        {
                            uint Temp = atomic_inc( &pQueue->m_controls.m_EnqueueMarkerScheduled );
                            if( Temp < MAX_NUMBER_OF_ENQUEUE_MARKER )
                            {
                                pCommand->m_commandState = CAN_BE_RECLAIMED;
                                commandsStack[ CurrentOffset ] = 0;
                                //Add events dependant on this command to list of events neeeded to be updated.
                                if( ( uint )pCommand->m_event != IGIL_EVENT_INVALID_HANDLE )
                                {
                                    pQueue->m_controls.m_EventDependencies[ atomic_inc( &pQueue->m_controls.m_CurrentScheduleEventNumber ) ] = pCommand->m_event;
                                }

                                //Remove event dependencies setting.
                                if( pCommand->m_numDependencies > 0 )
                                {
                                    DecreaseEventDependenciesParallel( pCommand, eventsPool );
                                }
                            }
                        }
                    }
                }
            }
            CurrentOffset += get_num_groups( 0 );
            if( pQueue->m_controls.m_IsSimulation )
            {
                barrier( CLK_LOCAL_MEM_FENCE );
            }
        }
    }

    //Finish execution and check end conditons
    //Execute this global barrier only when needed, i.e. stack browsing was executed or new item were added on the stack
    if( ( pQueue->m_controls.m_PreviousStackTop != pQueue->m_controls.m_StackSize ) |
        ( pQueue->m_controls.m_StackTop != pQueue->m_controls.m_PreviousStackTop ) )
    {
        GlobalBarrier( queueStorageBuffer );
    }

    //Cleanup & resource reclamation section
    //We are after global sync section, we can do anything to globals right now.
    if( ( get_local_id( 0 ) == 0 ) & ( get_group_id( 0 ) == 0 ) )
    {
        {
            pQueue->m_controls.m_CurrentDSHoffset           = pQueue->m_controls.m_DynamicHeapStart;
            pQueue->m_controls.m_IDTAfterFirstPhase         = 1;
            pQueue->m_controls.m_PreviousNumberOfQueues     = pQueue->m_controls.m_TotalNumberOfQueues;
            pQueue->m_controls.m_LastScheduleEventNumber    = pQueue->m_controls.m_CurrentScheduleEventNumber;
            pQueue->m_controls.m_CurrentScheduleEventNumber = 0;
        }
    }

    //Schedule scheduler
    if( ( (get_group_id( 0 ) == 1) == (get_num_groups( 0 ) > 1)  ) & ( get_local_id(0) == 0 ) )
    {
        pQueue->m_controls.m_SecondLevelBatchOffset = ( pQueue->m_controls.m_SecondLevelBatchOffset % MAX_SLB_OFFSET );

        //If we scheduled any blocks, put scheduler right after
        if( ( pQueue->m_controls.m_CurrentIDToffset > 1 ) | ( pQueue->m_controls.m_EnqueueMarkerScheduled > 0 ) )
        {
            AddCmdsInSLBforScheduler20Parallel( pQueue->m_controls.m_SecondLevelBatchOffset,
                                                pQueue,
                                                secondaryBatchBuffer,
                                                dsh );

            //If we have profiling enabled, we need CL_COMPLETE time, which is before next scheduler starts
            if( pQueue->m_controls.m_IsProfilingEnabled != 0 )
            {
                ulong Address  = ( ulong ) &( eventsPool->m_CLcompleteTimestamp );
                //Emit pipecontrol with timestamp write
                PatchPipeControlProfilingAddres( secondaryBatchBuffer,
                                                 pQueue->m_controls.m_SecondLevelBatchOffset,
                                                 Address,
                                                 PIPE_CONTROL_FOR_TIMESTAMP_START_OFFSET );
                //Bit after scheduler may be set by some other command, reset it to 0
                DisablePostSyncBitInPipeControl( secondaryBatchBuffer,
                                                 pQueue->m_controls.m_SecondLevelBatchOffset,
                                                 PIPE_CONTROL_FOR_TIMESTAMP_END_OFFSET_TO_PATCH );
#ifdef ENABLE_DEBUG_BUFFER
                if( ( DebugQueue != 0 ) & ( DebugQueue->m_flags == DDB_SCHEDULER_PROFILING ) )
                {
                    //Store Current scheduler START time
                    *((__global ulong * ) (&DebugQueue->m_data[ atomic_add( &DebugQueue->m_offset, 2 ) ] )) = EventsPool->m_CLcompleteTimestamp;

                    //Set address to store next scheduler's END time
                    PatchPipeControlProfilingAddres( secondaryBatchBuffer,
                                                        pQueue->m_controls.m_SecondLevelBatchOffset,
                                                        ( ulong )(&DebugQueue->m_data[ atomic_add( &DebugQueue->m_offset, 2 ) ]),
                                                        PIPE_CONTROL_FOR_TIMESTAMP_END_OFFSET );
                }
#endif
            }

            //Program pipe controls around scheduler to make sure it is not executed concurently to blocks
            else
            {
                //Locate previous pipe control
                int PreviousOffset = pQueue->m_controls.m_SecondLevelBatchOffset - SECOND_LEVEL_BUFFER_SPACE_FOR_EACH_ENQUEUE;
                //If offset is negative it means we are first command after chaining
                if( PreviousOffset < 0 )
                {
                    PreviousOffset += MAX_SLB_OFFSET;
                }
                //Tighten previous pipecontrol
                uint DwordOffset = PreviousOffset / DWORD_SIZE_IN_BYTES;

                PutCSStallPipeControl( secondaryBatchBuffer, DwordOffset, PIPE_CONTROL_FOR_TIMESTAMP_END_OFFSET );

                //Now put pipe control after scheduler
                DwordOffset = pQueue->m_controls.m_SecondLevelBatchOffset / DWORD_SIZE_IN_BYTES;

                PutCSStallPipeControl( secondaryBatchBuffer, DwordOffset, PIPE_CONTROL_FOR_TIMESTAMP_END_OFFSET );
            }

            pQueue->m_controls.m_SecondLevelBatchOffset  += SECOND_LEVEL_BUFFER_SPACE_FOR_EACH_ENQUEUE;
            pQueue->m_controls.m_CurrentIDToffset        = 1;
            pQueue->m_controls.m_EnqueueMarkerScheduled  = 0;
            pQueue->m_controls.Temporary[1]++;
            pQueue->m_controls.m_CurrentSSHoffset        = pQueue->m_controls.m_BTbaseOffset;
        }
        //Nothing to schedule, return to the host
        else
        {
#ifdef ENABLE_DEBUG_BUFFER
            //Set START time of current (last) scheduler
            if( ( pQueue->m_controls.m_IsProfilingEnabled != 0 ) & ( DebugQueue != 0 ) & ( DebugQueue->m_flags == DDB_SCHEDULER_PROFILING ) )
            {
                *((__global ulong * ) (&DebugQueue->m_data[ atomic_add( &DebugQueue->m_offset, 2 ) ] )) = EventsPool->m_CLcompleteTimestamp;
            }
#endif
            pQueue->m_controls.Temporary[2]++;
            pQueue->m_controls.m_SLBENDoffsetInBytes = ( int ) pQueue->m_controls.m_SecondLevelBatchOffset;
            //SlbOffset is expressed in bytes and for cmd it is needed to convert it to dwords
            __private uint DwordOffset = pQueue->m_controls.m_SecondLevelBatchOffset / DWORD_SIZE_IN_BYTES;
            //BB_START 1st DWORD
            secondaryBatchBuffer[ DwordOffset ] = OCLRT_BATCH_BUFFER_BEGIN_CMD_DWORD0;
            DwordOffset++;
            //BB_START 2nd DWORD - Address, 3rd DWORD Address high
            secondaryBatchBuffer[ DwordOffset++ ] = (uint)( pQueue->m_controls.m_CleanupSectionAddress & 0xFFFFFFFF );
            secondaryBatchBuffer[ DwordOffset++ ] = (uint)( ( pQueue->m_controls.m_CleanupSectionAddress >> 32 ) & 0xFFFFFFFF );
        }
    }

    //Parallel stack compaction
    if( ( ( get_group_id( 0 ) == 2 ) == ( get_num_groups( 0 ) > 2 )  ) & ( get_local_id( 0 ) == 0 ) )
    {
        uint Current = pQueue->m_controls.m_StackTop + get_local_id( 0 );
        uint StackSize = pQueue->m_controls.m_StackSize;
        uint Found = 0;

        while( ( Current < StackSize ) && ( Found == 0 ) )
        {
            __global uint * pCmdStackBlock = (__global uint *)( commandsStack + Current );
            //We have found an element
            if( *pCmdStackBlock != 0 )
            {
                Found = 1;
            }
            else
            {
                Current += get_local_size( 0 );
            }
        }

        if ( Found == 1 )
        {
            atomic_min( &pQueue->m_controls.m_StackTop, Current );
            atomic_min( &pQueue->m_controls.m_PreviousStackTop, Current );
        }
    }

    //Qstorage compaction
    if( ( ( get_group_id( 0 ) == 3 ) == ( get_num_groups( 0 ) > 3 )  ) & ( get_local_id( 0 ) == 0 ) )
    {
        uint ReclaimFurhter = 1;
        while( ( pQueue->m_controls.m_QstorageTop < pQueue->m_controls.m_QstorageSize ) & ( ReclaimFurhter == 1 ) )
        {
            pCommand = GetCommandHeaderFromStorage( ( __global uint* ) queueStorageBuffer, pQueue->m_controls.m_QstorageTop );
            if( pCommand->m_commandState == CAN_BE_RECLAIMED )
            {
                pQueue->m_controls.m_QstorageTop += pCommand->m_commandSize;
            }
            else
            {
                ReclaimFurhter = 0;
            }
        }

        pQueue->m_controls.m_PreviousStorageTop      = pQueue->m_controls.m_QstorageTop;

#ifndef DISABLE_RESOURCE_RECLAMATION
        //Reclaim space on queue_t, do this only if there is enough space
        //1 KB is used for global barrier, make sure this space will never be used.
        if( pQueue->m_controls.m_QstorageTop - 1024 > pQueue->m_size )
        {
            //In this case we can take full queue_t next time we enter scheduler, so reclaim full space on queue_t
            pQueue->m_head = IGIL_DEVICE_QUEUE_HEAD_INIT;
        }
#endif
        pQueue->m_controls.m_PreviousHead           = pQueue->m_head;
    }
}