intel/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2025-12-30 09:58:55 +08:00
Code Issues Packages Projects Releases Wiki Activity

Fix queue casts which cause spirv to die with assert

Browse Source 
Fixes: intel/opencl-clang#46

This fixes ocloc build of scheduler.cl which dies with assert from spirv:
  SPIRV-LLVM-Translator/lib/SPIRV/SPIRVWriter.cpp:685

Signed-off-by: Dmitry Rogozhkin <dmitry.v.rogozhkin@intel.com>

https://github.com/intel/compute-runtime/pull/143

Change-Id: I5c9d4a75321008802b9d77ef4a6bad3812e1b8b4
This commit is contained in:
Dmitry Rogozhkin
2019-03-13 05:18:09 -07:00
committed by sys_ocldev
parent 13f9f3a929
commit 7a7bb50328
1 changed files with 10 additions and 10 deletions
Download Patch File Download Diff File Expand all files Collapse all files

20
runtime/scheduler/scheduler.cl
View File

@@ -1,5 +1,5 @@
/*
* Copyright (C) 2017-2018 Intel Corporation
* Copyright (C) 2017-2019 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@@ -155,18 +155,18 @@ __global IGIL_KernelData* IGIL_GetKernelData( __global IGIL_KernelDataHeader* pK
}
inline __global IGIL_CommandHeader* TEMP_IGIL_GetCommandHeader( queue_t q, uint offset )
inline __global IGIL_CommandHeader* TEMP_IGIL_GetCommandHeader( __global IGIL_CommandQueue* q, uint offset )
{
__global uchar *pQueueRaw = __builtin_astype( q, __global uchar* );
__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( queue_t q, uint initialOffset, uint number )
__global IGIL_CommandHeader* TEMP_IGIL_GetNthCommandHeader( __global IGIL_CommandQueue* q, uint initialOffset, uint number )
{
__global uchar *pQueueRaw = __builtin_astype( q, __global uchar* );
__global uchar *pQueueRaw = (__global uchar *) q;
__global IGIL_CommandHeader* pCommand = ( __global IGIL_CommandHeader* )( pQueueRaw + initialOffset );
uint Offset = initialOffset;
@@ -182,9 +182,9 @@ __global IGIL_CommandHeader* TEMP_IGIL_GetNthCommandHeader( queue_t q, uint init
}
//Make sure enough command packets are in command queue before calling this function.
uint TEMP_IGIL_GetNthCommandHeaderOffset( queue_t q, uint initialOffset, uint number )
uint TEMP_IGIL_GetNthCommandHeaderOffset( __global IGIL_CommandQueue* q, uint initialOffset, uint number )
{
__global uchar *pQueueRaw = __builtin_astype( q, __global uchar* );
__global uchar *pQueueRaw = (__global uchar *) q;
__global IGIL_CommandHeader* pCommand = ( __global IGIL_CommandHeader* )( pQueueRaw + initialOffset );
uint Offset = initialOffset;
@@ -2878,13 +2878,13 @@ void SchedulerParallel20(
{
if( CurrentPacket == GroupID )
{
offset = TEMP_IGIL_GetNthCommandHeaderOffset( __builtin_astype( pQueue, queue_t ), InitialOffset, CurrentPacket );
offset = TEMP_IGIL_GetNthCommandHeaderOffset( pQueue, InitialOffset, CurrentPacket );
}
else
{
offset = TEMP_IGIL_GetNthCommandHeaderOffset( __builtin_astype( pQueue, queue_t ), offset, get_num_groups( 0 ) );
offset = TEMP_IGIL_GetNthCommandHeaderOffset( pQueue, offset, get_num_groups( 0 ) );
}
pCommand = TEMP_IGIL_GetCommandHeader( __builtin_astype( pQueue, queue_t ), offset );
pCommand = TEMP_IGIL_GetCommandHeader( pQueue, offset );
//Initialize command packet with proper lws
if( get_local_id( 0 ) == 0 )