Media/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2025-09-15 13:01:45 +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 8cc0177f1c
commit a39bc7e7b3
15 changed files with 202 additions and 49 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@ -1,5 +1,5 @@
/*
* Copyright (C) 2020-2021 Intel Corporation
* Copyright (C) 2020-2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@ -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(

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

@ -1,5 +1,5 @@
/*
* Copyright (C) 2020-2021 Intel Corporation
* Copyright (C) 2020-2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@ -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(

4
shared/test/common/helpers/kernel_binary_helper_hash_value.cpp
View File

@ -7,5 +7,5 @@
#include "shared/test/common/helpers/kernel_binary_helper.h"
const std::string KernelBinaryHelper::BUILT_INS("7998916142903730155");
const std::string KernelBinaryHelper::BUILT_INS_WITH_IMAGES("16526264370178379440_images");
const std::string KernelBinaryHelper::BUILT_INS("11256751929276190336");
const std::string KernelBinaryHelper::BUILT_INS_WITH_IMAGES("13967618058110882853_images");

20
shared/test/common/test_files/7998916142903730155.cl → shared/test/common/test_files/11256751929276190336.cl
View File

@ -1,5 +1,5 @@
/*
* Copyright (C) 2020-2022 Intel Corporation
* Copyright (C) 2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@ -156,13 +156,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(

0
shared/test/common/test_files/16526264370178379440_images_options.txt → shared/test/common/test_files/11256751929276190336_options.txt
View File

18
shared/test/common/test_files/16526264370178379440_images.cl → shared/test/common/test_files/13967618058110882853_images.cl
View File

@ -156,13 +156,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(

0
shared/test/common/test_files/7998916142903730155_options.txt → shared/test/common/test_files/13967618058110882853_images_options.txt
View File

14
shared/test/common/test_files/builtin_copyfill.cl
View File

@ -156,13 +156,21 @@ __kernel void FillBufferRightLeftover(
}
__kernel void FillBufferImmediate(
__global uint4* ptr,
uint 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 FillBufferImmediateRightLeftOver(
__global uchar* ptr,
uint 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 FillBufferSSHOffset(

14
shared/test/common/test_files/builtin_copyfill_stateless.cl
View File

@ -156,13 +156,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)
{
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)
{
size_t 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 FillBufferSSHOffset(

14
shared/test/common/test_files/builtin_images.cl
View File

@ -156,13 +156,21 @@ __kernel void FillBufferRightLeftover(
}
__kernel void FillBufferImmediate(
__global uint4* ptr,
uint 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 FillBufferImmediateRightLeftOver(
__global uchar* ptr,
uint 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 FillBufferSSHOffset(