feature: add built-in copy kernels for 1D_BUFFER images

Resolves: NEO-14782, HSD-18042093957

Signed-off-by: Igor Venevtsev <igor.venevtsev@intel.com>

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);
+}
 )==="

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");

shared/test/common/test_files/10604557797565794273_images_options.txt

@@ -1,8 +0,0 @@
-/*
- * Copyright (C) 2021-2025 Intel Corporation
- *
- * SPDX-License-Identifier: MIT
- *
- */
-
--cl-fast-relaxed-math

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,

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

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,