intel/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2025-12-25 05:24:02 +08:00
Code Issues Packages Projects Releases Wiki Activity

feature: add FillImage1dBuffer built-in kernel

Browse Source 
Resolves: NEO-13984, HSD-18041422852

Signed-off-by: Igor Venevtsev <igor.venevtsev@intel.com>
This commit is contained in:
Igor Venevtsev
2025-04-10 12:25:54 +00:00
committed by Compute-Runtime-Automation
parent ff7bfc8b6e
commit ef52479ce4
19 changed files with 1123 additions and 937 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
shared/source/built_ins/CMakeLists.txt
View File

@@ -1,5 +1,5 @@
#
# Copyright (C) 2020-2023 Intel Corporation
# Copyright (C) 2020-2025 Intel Corporation
#
# SPDX-License-Identifier: MIT
#
@@ -40,6 +40,7 @@ set(NEO_CORE_SRCS_BUILT_IN_KERNELS
${CMAKE_CURRENT_SOURCE_DIR}/kernels/fill_image1d.builtin_kernel
${CMAKE_CURRENT_SOURCE_DIR}/kernels/fill_image2d.builtin_kernel
${CMAKE_CURRENT_SOURCE_DIR}/kernels/fill_image3d.builtin_kernel
${CMAKE_CURRENT_SOURCE_DIR}/kernels/fill_image1d_buffer.builtin_kernel
)
set_property(GLOBAL PROPERTY NEO_CORE_SRCS_BUILT_IN_KERNELS ${NEO_CORE_SRCS_BUILT_IN_KERNELS})

9
shared/source/built_ins/built_in_ops_base.h
View File

@@ -1,5 +1,5 @@
/*
* Copyright (C) 2019-2024 Intel Corporation
* Copyright (C) 2019-2025 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@@ -43,6 +43,8 @@ inline constexpr Type fillImage2dHeapless{25};
inline constexpr Type fillImage3d{26};
inline constexpr Type fillImage3dHeapless{27};
inline constexpr Type queryKernelTimestamps{28};
inline constexpr Type fillImage1dBuffer{29};
inline constexpr Type fillImage1dBufferHeapless{30};
constexpr bool isStateless(Type type) {
constexpr std::array<Type, 10> statelessBuiltins{{copyBufferToBufferStateless, copyBufferRectStateless, fillBufferStateless, copyBufferToImage3dStateless,
@@ -59,7 +61,7 @@ constexpr bool isStateless(Type type) {
constexpr bool isHeapless(Type type) {
constexpr Type statelessBuiltins[] = {copyBufferToBufferStatelessHeapless, copyBufferRectStatelessHeapless, fillBufferStatelessHeapless,
copyBufferToImage3dHeapless, copyImage3dToBufferHeapless, copyImageToImage1dHeapless, copyImageToImage2dHeapless, copyImageToImage3dHeapless,
fillImage1dHeapless, fillImage2dHeapless, fillImage3dHeapless};
fillImage1dHeapless, fillImage2dHeapless, fillImage3dHeapless, fillImage1dBufferHeapless};
for (auto builtinType : statelessBuiltins) {
if (type == builtinType) {
@@ -141,8 +143,9 @@ DEFINE_ADJUST_BUILTIN_TYPE_IMAGE(copyImageToImage3d);
DEFINE_ADJUST_BUILTIN_TYPE_IMAGE(fillImage1d);
DEFINE_ADJUST_BUILTIN_TYPE_IMAGE(fillImage2d);
DEFINE_ADJUST_BUILTIN_TYPE_IMAGE(fillImage3d);
DEFINE_ADJUST_BUILTIN_TYPE_IMAGE(fillImage1dBuffer);
inline constexpr Type maxBaseValue{28};
inline constexpr Type maxBaseValue{fillImage1dBufferHeapless};
inline constexpr Type count{64};
} // namespace EBuiltInOps
} // namespace NEO

3
shared/source/built_ins/built_ins_storage.cpp
View File

@@ -73,6 +73,9 @@ const char *getBuiltinAsString(EBuiltInOps::Type builtin) {
return "fill_image3d.builtin_kernel";
case EBuiltInOps::queryKernelTimestamps:
return "copy_kernel_timestamps.builtin_kernel";
case EBuiltInOps::fillImage1dBuffer:
case EBuiltInOps::fillImage1dBufferHeapless:
return "fill_image1d_buffer.builtin_kernel";
};
}

1
shared/source/built_ins/builtins_binary.cmake
View File

@@ -31,6 +31,7 @@ set(GENERATED_BUILTINS_IMAGES
"fill_image1d"
"fill_image2d"
"fill_image3d"
"fill_image1d_buffer"
)
set(GENERATED_BUILTINS_IMAGES_STATELESS

18
shared/source/built_ins/kernels/fill_image1d_buffer.builtin_kernel Normal file
View File

@@ -0,0 +1,18 @@
/*
* Copyright (C) 2025 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
R"===(
__kernel void FillImage1dBuffer(
__write_only image1d_buffer_t output,
uint4 color,
int4 dstOffset) {
const int x = get_global_id(0);
const int dstCoord = x + dstOffset.x;
write_imageui(output, dstCoord, color);
}
)==="

13
shared/source/built_ins/registry/register_copy_kernels_source.cpp
View File

@@ -1,5 +1,5 @@
/*
* Copyright (C) 2020-2024 Intel Corporation
* Copyright (C) 2020-2025 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@@ -194,11 +194,20 @@ static RegisterEmbeddedResource registerAuxTranslationSrc(
static RegisterEmbeddedResource registerCopyKernelTimestampsSrc(
createBuiltinResourceName(
EBuiltInOps::fillImage3d,
EBuiltInOps::queryKernelTimestamps,
BuiltinCode::getExtension(BuiltinCode::ECodeType::source))
.c_str(),
std::string(
#include "shared/source/built_ins/kernels/copy_kernel_timestamps.builtin_kernel"
));
static RegisterEmbeddedResource registerFillImage1dBufferSrc(
createBuiltinResourceName(
EBuiltInOps::fillImage1dBuffer,
BuiltinCode::getExtension(BuiltinCode::ECodeType::source))
.c_str(),
std::string(
#include "shared/source/built_ins/kernels/fill_image1d_buffer.builtin_kernel"
));
} // namespace NEO