From 42efb3d2047381e39a4ca51f5c23db66b2157bc8 Mon Sep 17 00:00:00 2001 From: Igor Venevtsev Date: Thu, 1 May 2025 00:48:10 +0000 Subject: [PATCH] feature: add built-in copy kernels for 1D_BUFFER images Resolves: NEO-14782, HSD-18042093957 Signed-off-by: Igor Venevtsev --- .../builtin/builtin_functions_lib_impl.cpp | 6 +-- .../builtin/builtin_functions_tests.cpp | 4 +- manifests/manifest.yml | 2 +- .../built_ins/builtins_dispatch_builder.cpp | 33 +++++++++++-- .../enqueue_copy_image_tests.cpp | 48 ++++++++++++++++++ .../get_size_required_image_tests.cpp | 4 +- .../copy_image_to_image3d.builtin_kernel | 49 ++++++++++++++++++- .../kernel_binary_helper_hash_value.cpp | 2 +- .../10604557797565794273_images_options.txt | 8 --- ...mages.cl => 2205520382307710565_images.cl} | 47 +++++++++++++++++- .../2205520382307710565_images_options.txt | 8 +++ .../test/common/test_files/builtin_images.cl | 47 +++++++++++++++++- 12 files changed, 232 insertions(+), 26 deletions(-) delete mode 100644 shared/test/common/test_files/10604557797565794273_images_options.txt rename shared/test/common/test_files/{10604557797565794273_images.cl => 2205520382307710565_images.cl} (95%) create mode 100644 shared/test/common/test_files/2205520382307710565_images_options.txt diff --git a/level_zero/core/source/builtin/builtin_functions_lib_impl.cpp b/level_zero/core/source/builtin/builtin_functions_lib_impl.cpp index 7f814939e0..6d6f29911a 100644 --- a/level_zero/core/source/builtin/builtin_functions_lib_impl.cpp +++ b/level_zero/core/source/builtin/builtin_functions_lib_impl.cpp @@ -1,5 +1,5 @@ /* - * Copyright (C) 2020-2024 Intel Corporation + * Copyright (C) 2020-2025 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -285,11 +285,11 @@ void BuiltinFunctionsLibImpl::initBuiltinImageKernel(ImageBuiltin func) { builtin = NEO::EBuiltInOps::copyImage3dToBufferHeapless; break; case ImageBuiltin::copyImageRegion: - builtinName = "CopyImageToImage3d"; + builtinName = "CopyImage3dToImage3d"; builtin = NEO::EBuiltInOps::copyImageToImage3d; break; case ImageBuiltin::copyImageRegionHeapless: - builtinName = "CopyImageToImage3d"; + builtinName = "CopyImage3dToImage3d"; builtin = NEO::EBuiltInOps::copyImageToImage3dHeapless; break; default: diff --git a/level_zero/core/test/unit_tests/sources/builtin/builtin_functions_tests.cpp b/level_zero/core/test/unit_tests/sources/builtin/builtin_functions_tests.cpp index 377ed9ffd3..d18ff9e8ff 100644 --- a/level_zero/core/test/unit_tests/sources/builtin/builtin_functions_tests.cpp +++ b/level_zero/core/test/unit_tests/sources/builtin/builtin_functions_tests.cpp @@ -1,5 +1,5 @@ /* - * Copyright (C) 2020-2024 Intel Corporation + * Copyright (C) 2020-2025 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -238,7 +238,7 @@ HWTEST_F(TestBuiltinFunctionsLibImpl, givenHeaplessImageBuiltinsWhenInitBuiltinK lib.initBuiltinImageKernel(L0::ImageBuiltin::copyImageRegionHeapless); EXPECT_EQ(NEO::EBuiltInOps::copyImageToImage3dHeapless, lib.builtinPassed); - EXPECT_STREQ("CopyImageToImage3d", lib.kernelNamePassed.c_str()); + EXPECT_STREQ("CopyImage3dToImage3d", lib.kernelNamePassed.c_str()); } HWTEST_F(TestBuiltinFunctionsLibImpl, givenCompilerInterfaceWhenCreateDeviceAndImageSupportedThenBuiltinsImageFunctionsAreLoaded) { diff --git a/manifests/manifest.yml b/manifests/manifest.yml index 000b70b8ba..9b426e9356 100644 --- a/manifests/manifest.yml +++ b/manifests/manifest.yml @@ -41,7 +41,7 @@ components: dest_dir: kernels_bin type: git branch: kernels_bin - revision: 3394-4091 + revision: 3394-4097 kmdaf: branch: kmdaf dest_dir: kmdaf diff --git a/opencl/source/built_ins/builtins_dispatch_builder.cpp b/opencl/source/built_ins/builtins_dispatch_builder.cpp index e9ec52d5d3..16fd692c42 100644 --- a/opencl/source/built_ins/builtins_dispatch_builder.cpp +++ b/opencl/source/built_ins/builtins_dispatch_builder.cpp @@ -856,7 +856,10 @@ class BuiltInOp : public BuiltinDispatchInfoBui : BuiltinDispatchInfoBuilder(kernelsLib, device) { populate(EBuiltInOps::copyImageToImage3d, "", - "CopyImageToImage3d", kernel); + "CopyImage3dToImage3d", kernelCopyImage3dToImage3d, + "CopyImage1dBufferToImage3d", kernelCopyImage1dBufferToImage3d, + "CopyImage3dToImage1dBuffer", kernelCopyImage3dToImage1dBuffer, + "CopyImage1dBufferToImage1dBuffer", kernelCopyImage1dBufferToImage1dBuffer); } bool buildDispatchInfos(MultiDispatchInfo &multiDispatchInfo) const override { @@ -874,8 +877,19 @@ class BuiltInOp : public BuiltinDispatchInfoBui multiDispatchInfo.pushRedescribedMemObj(std::unique_ptr(srcImageRedescribed)); // life range same as mdi's multiDispatchInfo.pushRedescribedMemObj(std::unique_ptr(dstImageRedescribed)); // life range same as mdi's + bool src1dBuffer = srcImage->getImageDesc().image_type == CL_MEM_OBJECT_IMAGE1D_BUFFER; + bool dst1dBuffer = dstImage->getImageDesc().image_type == CL_MEM_OBJECT_IMAGE1D_BUFFER; + // Set-up kernel - kernelNoSplit3DBuilder.setKernel(kernel->getKernel(clDevice.getRootDeviceIndex())); + if (src1dBuffer && dst1dBuffer) { + kernelNoSplit3DBuilder.setKernel(kernelCopyImage1dBufferToImage1dBuffer->getKernel(clDevice.getRootDeviceIndex())); + } else if (src1dBuffer) { + kernelNoSplit3DBuilder.setKernel(kernelCopyImage1dBufferToImage3d->getKernel(clDevice.getRootDeviceIndex())); + } else if (dst1dBuffer) { + kernelNoSplit3DBuilder.setKernel(kernelCopyImage3dToImage1dBuffer->getKernel(clDevice.getRootDeviceIndex())); + } else { + kernelNoSplit3DBuilder.setKernel(kernelCopyImage3dToImage3d->getKernel(clDevice.getRootDeviceIndex())); + } // Set-up source image kernelNoSplit3DBuilder.setArg(0, srcImageRedescribed, operationParams.srcMipLevel); @@ -916,11 +930,17 @@ class BuiltInOp : public BuiltinDispatchInfoBui if (populateKernels) { populate(EBuiltInOps::copyImageToImage3d, "", - "CopyImageToImage3d", kernel); + "CopyImage3dToImage3d", kernelCopyImage3dToImage3d, + "CopyImage1dBufferToImage3d", kernelCopyImage1dBufferToImage3d, + "CopyImage3dToImage1dBuffer", kernelCopyImage3dToImage1dBuffer, + "CopyImage1dBufferToImage1dBuffer", kernelCopyImage1dBufferToImage1dBuffer); } } - MultiDeviceKernel *kernel = nullptr; + MultiDeviceKernel *kernelCopyImage3dToImage3d = nullptr; + MultiDeviceKernel *kernelCopyImage1dBufferToImage3d = nullptr; + MultiDeviceKernel *kernelCopyImage3dToImage1dBuffer = nullptr; + MultiDeviceKernel *kernelCopyImage1dBufferToImage1dBuffer = nullptr; }; template <> @@ -930,7 +950,10 @@ class BuiltInOp : public BuiltInOp(kernelsLib, device, false) { populate(EBuiltInOps::copyImageToImage3dHeapless, "", - "CopyImageToImage3d", kernel); + "CopyImage3dToImage3d", kernelCopyImage3dToImage3d, + "CopyImage1dBufferToImage3d", kernelCopyImage1dBufferToImage3d, + "CopyImage3dToImage1dBuffer", kernelCopyImage3dToImage1dBuffer, + "CopyImage1dBufferToImage1dBuffer", kernelCopyImage1dBufferToImage1dBuffer); } }; diff --git a/opencl/test/unit_test/command_queue/enqueue_copy_image_tests.cpp b/opencl/test/unit_test/command_queue/enqueue_copy_image_tests.cpp index 4a4bc20443..5c7ac02650 100644 --- a/opencl/test/unit_test/command_queue/enqueue_copy_image_tests.cpp +++ b/opencl/test/unit_test/command_queue/enqueue_copy_image_tests.cpp @@ -380,3 +380,51 @@ HWTEST_F(OneMipLevelCopyImageImageTests, GivenNotMippedImageWhenCopyingImageThen EXPECT_EQ(0u, usedBuiltinOpsParams.srcMipLevel); EXPECT_EQ(0u, usedBuiltinOpsParams.dstMipLevel); } + +HWTEST_F(EnqueueCopyImageTest, WhenCopyImage1dBufferToImage1dBufferThenCorrectBuitInIsSelected) { + auto mockCmdQ = std::make_unique>(context, pClDevice, nullptr); + VariableBackup cmdQBackup(&pCmdQ, mockCmdQ.get()); + std::unique_ptr srcImage1dBuffer; + srcImage1dBuffer.reset(Image1dBufferHelper<>::create(context)); + VariableBackup srcImageBackup(&srcImage, srcImage1dBuffer.get()); + std::unique_ptr dstImage1dBuffer; + dstImage1dBuffer.reset(Image1dBufferHelper<>::create(context)); + VariableBackup dstImageBackup(&dstImage, dstImage1dBuffer.get()); + mockCmdQ->storeMultiDispatchInfo = true; + EnqueueCopyImageHelper<>::enqueueCopyImage(pCmdQ, srcImage, dstImage); + const auto &kernelInfo = mockCmdQ->storedMultiDispatchInfo.begin()->getKernel()->getKernelInfo(); + EXPECT_TRUE(kernelInfo.kernelDescriptor.kernelMetadata.kernelName == "CopyImage1dBufferToImage1dBuffer"); +} + +HWTEST_F(EnqueueCopyImageTest, WhenCopyImage1dBufferToImageThenCorrectBuitInIsSelected) { + auto mockCmdQ = std::make_unique>(context, pClDevice, nullptr); + VariableBackup cmdQBackup(&pCmdQ, mockCmdQ.get()); + std::unique_ptr srcImage1dBuffer; + srcImage1dBuffer.reset(Image1dBufferHelper<>::create(context)); + VariableBackup srcImageBackup(&srcImage, srcImage1dBuffer.get()); + mockCmdQ->storeMultiDispatchInfo = true; + EnqueueCopyImageHelper<>::enqueueCopyImage(pCmdQ, srcImage, dstImage); + const auto &kernelInfo = mockCmdQ->storedMultiDispatchInfo.begin()->getKernel()->getKernelInfo(); + EXPECT_TRUE(kernelInfo.kernelDescriptor.kernelMetadata.kernelName == "CopyImage1dBufferToImage3d"); +} + +HWTEST_F(EnqueueCopyImageTest, WhenCopyImageToImage1dBufferThenCorrectBuitInIsSelected) { + auto mockCmdQ = std::make_unique>(context, pClDevice, nullptr); + VariableBackup cmdQBackup(&pCmdQ, mockCmdQ.get()); + std::unique_ptr dstImage1dBuffer; + dstImage1dBuffer.reset(Image1dBufferHelper<>::create(context)); + VariableBackup dstImageBackup(&dstImage, dstImage1dBuffer.get()); + mockCmdQ->storeMultiDispatchInfo = true; + EnqueueCopyImageHelper<>::enqueueCopyImage(pCmdQ, srcImage, dstImage); + const auto &kernelInfo = mockCmdQ->storedMultiDispatchInfo.begin()->getKernel()->getKernelInfo(); + EXPECT_TRUE(kernelInfo.kernelDescriptor.kernelMetadata.kernelName == "CopyImage3dToImage1dBuffer"); +} + +HWTEST_F(EnqueueCopyImageTest, WhenCopyImageToImageThenCorrectBuitInIsSelected) { + auto mockCmdQ = std::make_unique>(context, pClDevice, nullptr); + VariableBackup cmdQBackup(&pCmdQ, mockCmdQ.get()); + mockCmdQ->storeMultiDispatchInfo = true; + EnqueueCopyImageHelper<>::enqueueCopyImage(pCmdQ, srcImage, dstImage); + const auto &kernelInfo = mockCmdQ->storedMultiDispatchInfo.begin()->getKernel()->getKernelInfo(); + EXPECT_TRUE(kernelInfo.kernelDescriptor.kernelMetadata.kernelName == "CopyImage3dToImage3d"); +} diff --git a/opencl/test/unit_test/command_queue/get_size_required_image_tests.cpp b/opencl/test/unit_test/command_queue/get_size_required_image_tests.cpp index 3a2b770772..5a54a7ea3d 100644 --- a/opencl/test/unit_test/command_queue/get_size_required_image_tests.cpp +++ b/opencl/test/unit_test/command_queue/get_size_required_image_tests.cpp @@ -124,10 +124,10 @@ HWTEST_F(GetSizeRequiredImageTest, WhenCopyingReadWriteImageThenHeapsAndCommandB auto usedBeforeIOH = ioh.getUsed(); auto usedBeforeSSH = ssh.getUsed(); - std::unique_ptr program(Program::createBuiltInFromSource("CopyImageToImage3d", context, context->getDevices(), nullptr)); + std::unique_ptr program(Program::createBuiltInFromSource("CopyImageTo3dImage3d", context, context->getDevices(), nullptr)); program->build(program->getDevices(), nullptr); cl_int retVal{CL_SUCCESS}; - std::unique_ptr kernel(Kernel::create(program.get(), program->getKernelInfoForKernel("CopyImageToImage3d"), *context->getDevice(0), retVal)); + std::unique_ptr kernel(Kernel::create(program.get(), program->getKernelInfoForKernel("CopyImage3dToImage3d"), *context->getDevice(0), retVal)); EXPECT_EQ(CL_SUCCESS, retVal); EXPECT_NE(nullptr, kernel); diff --git a/shared/source/built_ins/kernels/copy_image_to_image3d.builtin_kernel b/shared/source/built_ins/kernels/copy_image_to_image3d.builtin_kernel index fc35229c70..dffd480b67 100644 --- a/shared/source/built_ins/kernels/copy_image_to_image3d.builtin_kernel +++ b/shared/source/built_ins/kernels/copy_image_to_image3d.builtin_kernel @@ -1,5 +1,5 @@ /* - * Copyright (C) 2018-2021 Intel Corporation + * Copyright (C) 2018-2025 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -8,7 +8,7 @@ R"===( #pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable -__kernel void CopyImageToImage3d( +__kernel void CopyImage3dToImage3d( __read_only image3d_t input, __write_only image3d_t output, int4 srcOffset, @@ -22,4 +22,49 @@ __kernel void CopyImageToImage3d( const uint4 c = read_imageui(input, srcCoord); write_imageui(output, dstCoord, c); } + +__kernel void CopyImage1dBufferToImage3d( + __read_only image1d_buffer_t input, + __write_only image3d_t output, + int4 srcOffset, + int4 dstOffset) { + const int x = get_global_id(0); + const int y = get_global_id(1); + const int z = get_global_id(2); + + const int4 srcCoord = (int4)(x, y, z, 0) + srcOffset; + const int4 dstCoord = (int4)(x, y, z, 0) + dstOffset; + const uint4 c = read_imageui(input, srcCoord.x); + write_imageui(output, dstCoord, c); +} + +__kernel void CopyImage3dToImage1dBuffer( + __read_only image3d_t input, + __write_only image1d_buffer_t output, + int4 srcOffset, + int4 dstOffset) { + const int x = get_global_id(0); + const int y = get_global_id(1); + const int z = get_global_id(2); + + const int4 srcCoord = (int4)(x, y, z, 0) + srcOffset; + const int4 dstCoord = (int4)(x, y, z, 0) + dstOffset; + const uint4 c = read_imageui(input, srcCoord); + write_imageui(output, dstCoord.x, c); +} + +__kernel void CopyImage1dBufferToImage1dBuffer( + __read_only image1d_buffer_t input, + __write_only image1d_buffer_t output, + int4 srcOffset, + int4 dstOffset) { + const int x = get_global_id(0); + const int y = get_global_id(1); + const int z = get_global_id(2); + + const int4 srcCoord = (int4)(x, y, z, 0) + srcOffset; + const int4 dstCoord = (int4)(x, y, z, 0) + dstOffset; + const uint4 c = read_imageui(input, srcCoord.x); + write_imageui(output, dstCoord.x, c); +} )===" diff --git a/shared/test/common/helpers/kernel_binary_helper_hash_value.cpp b/shared/test/common/helpers/kernel_binary_helper_hash_value.cpp index d8259eac5d..07de709823 100644 --- a/shared/test/common/helpers/kernel_binary_helper_hash_value.cpp +++ b/shared/test/common/helpers/kernel_binary_helper_hash_value.cpp @@ -8,4 +8,4 @@ #include "shared/test/common/helpers/kernel_binary_helper.h" const std::string KernelBinaryHelper::BUILT_INS("15672580764041246108"); -const std::string KernelBinaryHelper::BUILT_INS_WITH_IMAGES("10604557797565794273_images"); +const std::string KernelBinaryHelper::BUILT_INS_WITH_IMAGES("2205520382307710565_images"); diff --git a/shared/test/common/test_files/10604557797565794273_images_options.txt b/shared/test/common/test_files/10604557797565794273_images_options.txt deleted file mode 100644 index 973f658574..0000000000 --- a/shared/test/common/test_files/10604557797565794273_images_options.txt +++ /dev/null @@ -1,8 +0,0 @@ -/* - * Copyright (C) 2021-2025 Intel Corporation - * - * SPDX-License-Identifier: MIT - * - */ - --cl-fast-relaxed-math diff --git a/shared/test/common/test_files/10604557797565794273_images.cl b/shared/test/common/test_files/2205520382307710565_images.cl similarity index 95% rename from shared/test/common/test_files/10604557797565794273_images.cl rename to shared/test/common/test_files/2205520382307710565_images.cl index 9324ef2658..fa064e56da 100644 --- a/shared/test/common/test_files/10604557797565794273_images.cl +++ b/shared/test/common/test_files/2205520382307710565_images.cl @@ -473,7 +473,7 @@ __kernel void CopyImageToImage2d( #pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable -__kernel void CopyImageToImage3d( +__kernel void CopyImage3dToImage3d( __read_only image3d_t input, __write_only image3d_t output, int4 srcOffset, @@ -488,6 +488,51 @@ __kernel void CopyImageToImage3d( write_imageui(output, dstCoord, c); } +__kernel void CopyImage1dBufferToImage3d( + __read_only image1d_buffer_t input, + __write_only image3d_t output, + int4 srcOffset, + int4 dstOffset) { + const int x = get_global_id(0); + const int y = get_global_id(1); + const int z = get_global_id(2); + + const int4 srcCoord = (int4)(x, y, z, 0) + srcOffset; + const int4 dstCoord = (int4)(x, y, z, 0) + dstOffset; + const uint4 c = read_imageui(input, srcCoord.x); + write_imageui(output, dstCoord, c); +} + +__kernel void CopyImage3dToImage1dBuffer( + __read_only image3d_t input, + __write_only image1d_buffer_t output, + int4 srcOffset, + int4 dstOffset) { + const int x = get_global_id(0); + const int y = get_global_id(1); + const int z = get_global_id(2); + + const int4 srcCoord = (int4)(x, y, z, 0) + srcOffset; + const int4 dstCoord = (int4)(x, y, z, 0) + dstOffset; + const uint4 c = read_imageui(input, srcCoord); + write_imageui(output, dstCoord.x, c); +} + +__kernel void CopyImage1dBufferToImage1dBuffer( + __read_only image1d_buffer_t input, + __write_only image1d_buffer_t output, + int4 srcOffset, + int4 dstOffset) { + const int x = get_global_id(0); + const int y = get_global_id(1); + const int z = get_global_id(2); + + const int4 srcCoord = (int4)(x, y, z, 0) + srcOffset; + const int4 dstCoord = (int4)(x, y, z, 0) + dstOffset; + const uint4 c = read_imageui(input, srcCoord.x); + write_imageui(output, dstCoord.x, c); +} + #pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable __kernel void CopyBufferToImage3dBytes(__global uchar *src, diff --git a/shared/test/common/test_files/2205520382307710565_images_options.txt b/shared/test/common/test_files/2205520382307710565_images_options.txt new file mode 100644 index 0000000000..40469c4dea --- /dev/null +++ b/shared/test/common/test_files/2205520382307710565_images_options.txt @@ -0,0 +1,8 @@ +/* + * Copyright (C) 2025 Intel Corporation + * + * SPDX-License-Identifier: MIT + * + */ + +-cl-fast-relaxed-math \ No newline at end of file diff --git a/shared/test/common/test_files/builtin_images.cl b/shared/test/common/test_files/builtin_images.cl index a444d11f94..b6da2e01db 100644 --- a/shared/test/common/test_files/builtin_images.cl +++ b/shared/test/common/test_files/builtin_images.cl @@ -428,7 +428,7 @@ __kernel void CopyImageToImage2d( #pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable -__kernel void CopyImageToImage3d( +__kernel void CopyImage3dToImage3d( __read_only image3d_t input, __write_only image3d_t output, int4 srcOffset, @@ -443,6 +443,51 @@ __kernel void CopyImageToImage3d( write_imageui(output, dstCoord, c); } +__kernel void CopyImage1dBufferToImage3d( + __read_only image1d_buffer_t input, + __write_only image3d_t output, + int4 srcOffset, + int4 dstOffset) { + const int x = get_global_id(0); + const int y = get_global_id(1); + const int z = get_global_id(2); + + const int4 srcCoord = (int4)(x, y, z, 0) + srcOffset; + const int4 dstCoord = (int4)(x, y, z, 0) + dstOffset; + const uint4 c = read_imageui(input, srcCoord.x); + write_imageui(output, dstCoord, c); +} + +__kernel void CopyImage3dToImage1dBuffer( + __read_only image3d_t input, + __write_only image1d_buffer_t output, + int4 srcOffset, + int4 dstOffset) { + const int x = get_global_id(0); + const int y = get_global_id(1); + const int z = get_global_id(2); + + const int4 srcCoord = (int4)(x, y, z, 0) + srcOffset; + const int4 dstCoord = (int4)(x, y, z, 0) + dstOffset; + const uint4 c = read_imageui(input, srcCoord); + write_imageui(output, dstCoord.x, c); +} + +__kernel void CopyImage1dBufferToImage1dBuffer( + __read_only image1d_buffer_t input, + __write_only image1d_buffer_t output, + int4 srcOffset, + int4 dstOffset) { + const int x = get_global_id(0); + const int y = get_global_id(1); + const int z = get_global_id(2); + + const int4 srcCoord = (int4)(x, y, z, 0) + srcOffset; + const int4 dstCoord = (int4)(x, y, z, 0) + dstOffset; + const uint4 c = read_imageui(input, srcCoord.x); + write_imageui(output, dstCoord.x, c); +} + #pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable __kernel void CopyBufferToImage3dBytes(__global uchar *src,