intel/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2026-01-04 23:56:39 +08:00
Code Issues Packages Projects Releases Wiki Activity

Use uint4 type instead of char in builtin kernel

Browse Source 
Related-To: NEO-7287

Signed-off-by: Szymon Morek <szymon.morek@intel.com>
This commit is contained in:
Szymon Morek
2022-08-26 12:56:55 +00:00
committed by Compute-Runtime-Automation
parent 2f9d0ad584
commit 25f4b06258
15 changed files with 218 additions and 47 deletions
Download Patch File Download Diff File Expand all files Collapse all files

18
shared/source/built_ins/kernels/fill_buffer.builtin_kernel
View File

@@ -48,13 +48,21 @@ __kernel void FillBufferRightLeftover(
}
__kernel void FillBufferImmediate(
__global uchar* ptr,
uint dstSshOffset, // Offset needed in case ptr has been adjusted for SSH alignment
__global uint4* ptr,
ulong dstSshOffset, // Offset needed in case ptr has been adjusted for SSH alignment
const uint value)
{
uint dstIndex = get_global_id(0);
__global uchar* pDst = (__global uchar*)ptr + dstSshOffset;
pDst[dstIndex] = value;
uint gid = get_global_id(0);
(ptr + dstSshOffset)[gid] = value;
}
__kernel void FillBufferImmediateRightLeftOver(
__global uchar* ptr,
ulong dstSshOffset, // Offset needed in case ptr has been adjusted for SSH alignment
const uint value)
{
uint gid = get_global_id(0);
(ptr + dstSshOffset)[gid] = value;
}
__kernel void FillBufferSSHOffset(

14
shared/source/built_ins/kernels/fill_buffer_stateless.builtin_kernel
View File

@@ -48,13 +48,21 @@ __kernel void FillBufferRightLeftover(
}
__kernel void FillBufferImmediate(
__global uint4* ptr,
ulong dstSshOffset, // Offset needed in case ptr has been adjusted for SSH alignment
const uint value)
{
size_t gid = get_global_id(0);
(ptr + dstSshOffset)[gid] = value;
}
__kernel void FillBufferImmediateRightLeftOver(
__global uchar* ptr,
ulong dstSshOffset, // Offset needed in case ptr has been adjusted for SSH alignment
const uint value)
{
size_t dstIndex = get_global_id(0);
__global uchar* pDst = (__global uchar*)ptr + dstSshOffset;
pDst[dstIndex] = value;
size_t gid = get_global_id(0);
(ptr + dstSshOffset)[gid] = value;
}
__kernel void FillBufferSSHOffset(