Aub test for bindless image

Related-To: NEO-4607

Change-Id: I0abef7d3b12f67eed7b5a11a3b9655f42de5a2cf
Signed-off-by: Mateusz Hoppe <mateusz.hoppe@intel.com>

opencl/source/kernel/kernel.cpp

@@ -2449,7 +2449,7 @@ uint64_t Kernel::getKernelStartOffset(
 void Kernel::patchBindlessSurfaceStateOffsets(const size_t sshOffset) {
     const bool bindlessBuffers = DebugManager.flags.UseBindlessBuffers.get();
     const bool bindlessImages = DebugManager.flags.UseBindlessImages.get();
-    const bool bindlessUsed = bindlessBuffers || bindlessImages;
+    const bool bindlessUsed = (bindlessBuffers || bindlessImages) && !isBuiltIn;
 
     if (bindlessUsed) {
         auto &hardwareInfo = getDevice().getHardwareInfo();

opencl/test/unit_test/CMakeLists.txt

@@ -261,9 +261,10 @@ function(neo_gen_kernels_with_options platform_name_with_type platform_name suff
   set(kernels_to_compile_${platform_name_with_type} ${kernels_to_compile_${platform_name_with_type}} PARENT_SCOPE)
 endfunction()
 
-function(neo_gen_kernels_with_internal_options platform_name_with_type platform_name suffix filepath output_name)
+function(neo_gen_kernels_with_internal_options platform_name_with_type platform_name suffix filepath output_name_prefix)
   set(kernels_to_compile)
-  foreach(filearg ${filepath})
+
+  set(filearg ${filepath})
   get_filename_component(filename ${filearg} NAME)
   get_filename_component(basename ${filearg} NAME_WE)
   get_filename_component(base_workdir ${filearg} DIRECTORY)
@@ -271,8 +272,8 @@ function(neo_gen_kernels_with_internal_options platform_name_with_type platform_
   set(outputdir "${TargetDir}/${suffix}/test_files/${NEO_ARCH}/")
   set(workdir "${CMAKE_CURRENT_SOURCE_DIR}/${base_workdir}/")
 
-    if (NOT "${output_name}" STREQUAL "")
-      set(basename ${output_name})
+  if (NOT "${output_name_prefix}" STREQUAL "")
+    set(basename ${output_name_prefix}_${basename})
   endif()
 
   set(outputpath_base "${outputdir}/${basename}_${suffix}")
@@ -282,19 +283,18 @@ function(neo_gen_kernels_with_internal_options platform_name_with_type platform_
     ${outputpath_base}.gen
   )
 
-    if (NOT "${output_name}" STREQUAL "")
-      set(output_name -output ${output_name})
+  if (NOT "${output_name_prefix}" STREQUAL "")
+    set(output_name -output ${basename})
   endif()
 
   add_custom_command(
     OUTPUT ${output_files}
-      COMMAND ${cloc_cmd_prefix} -q -file ${filename} -device ${platform_name} -${NEO_BITS} -out_dir ${outputdir} ${output_name} -internal_options ${ARGN}
+    COMMAND ${cloc_cmd_prefix}  -file ${filename} -device ${platform_name} -${NEO_BITS} -out_dir ${outputdir} ${output_name} -internal_options ${ARGN}
     WORKING_DIRECTORY ${workdir}
     DEPENDS ${filearg} ocloc
   )
 
   list(APPEND kernels_to_compile ${output_files})
-  endforeach()
   list(APPEND kernels_to_compile_${platform_name_with_type} ${kernels_to_compile})
   set(kernels_to_compile_${platform_name_with_type} ${kernels_to_compile_${platform_name_with_type}} PARENT_SCOPE)
 endfunction()
@@ -407,6 +407,7 @@ set(TEST_KERNEL_BINDLESS_internal_options
 
 set(TEST_KERNEL_BINDLESS
   test_files/stateful_copy_buffer.cl
+  test_files/copy_buffer_to_image.cl
 )
 
 file(GLOB_RECURSE TEST_KERNELS test_files/*.cl)
@@ -448,7 +449,9 @@ macro(macro_for_each_gen)
 
         # Temporarily disable in Debug builds
         if(NOT ${CMAKE_BUILD_TYPE} STREQUAL "Debug")
-          neo_gen_kernels_with_internal_options(${family_name_with_type} ${PLATFORM_LOWER} ${family_name_with_type} ${TEST_KERNEL_BINDLESS} "bindless_copy_buffer" ${TEST_KERNEL_BINDLESS_internal_options})
+          foreach(file ${TEST_KERNEL_BINDLESS})
+            neo_gen_kernels_with_internal_options(${family_name_with_type} ${PLATFORM_LOWER} ${family_name_with_type} ${file} "bindless" ${TEST_KERNEL_BINDLESS_internal_options})
+          endforeach()
         endif()
 
         set(sip_kernel_file_name)

opencl/test/unit_test/aub_tests/command_queue/enqueue_kernel_aub_tests.cpp

@@ -876,6 +876,9 @@ using IsSklPlus = IsAtLeastProduct<IGFX_SKYLAKE>;
 
 HWTEST2_F(AUBBindlessKernel, givenBindlessCopyKernelWhenEnqueuedThenResultsValidate, IsSklPlus) {
     constexpr size_t bufferSize = MemoryConstants::pageSize;
+
+    createKernel(std::string("bindless_stateful_copy_buffer"), std::string("StatefulCopyBuffer"));
+
     cl_uint workDim = 1;
     size_t globalWorkOffset[3] = {0, 0, 0};
     size_t globalWorkSize[3] = {bufferSize / 2, 1, 1};
@@ -948,3 +951,100 @@ HWTEST2_F(AUBBindlessKernel, givenBindlessCopyKernelWhenEnqueuedThenResultsValid
     expectMemory<FamilyType>(reinterpret_cast<void *>(pBufferDst->getGraphicsAllocation()->getGpuAddress()),
                              bufferDataSrc, bufferSize);
 }
+
+HWTEST2_F(AUBBindlessKernel, DISABLED_givenBindlessCopyImageKernelWhenEnqueuedThenResultsValidate, IsSklPlus) {
+    constexpr unsigned int testWidth = 5;
+    constexpr unsigned int testHeight = 1;
+    constexpr unsigned int testDepth = 1;
+
+    createKernel(std::string("bindless_copy_buffer_to_image"), std::string("CopyBufferToImage3d"));
+
+    constexpr size_t imageSize = testWidth * testHeight * testDepth;
+    cl_uint workDim = 1;
+    size_t globalWorkOffset[3] = {0, 0, 0};
+    size_t globalWorkSize[3] = {imageSize, 1, 1};
+    size_t localWorkSize[3] = {1, 1, 1};
+    cl_uint numEventsInWaitList = 0;
+    cl_event *eventWaitList = nullptr;
+    cl_event *event = nullptr;
+
+    uint8_t imageDataSrc[imageSize];
+    uint8_t imageDataDst[imageSize + 1];
+
+    memset(imageDataSrc, 1, imageSize);
+    memset(imageDataDst, 0, imageSize + 1);
+
+    cl_image_format imageFormat = {0};
+    cl_image_desc imageDesc = {0};
+
+    imageFormat.image_channel_data_type = CL_UNSIGNED_INT8;
+    imageFormat.image_channel_order = CL_R;
+    imageDesc.image_type = CL_MEM_OBJECT_IMAGE1D;
+    imageDesc.image_width = testWidth;
+    imageDesc.image_height = testHeight;
+    imageDesc.image_depth = testDepth;
+    imageDesc.image_array_size = 1;
+    imageDesc.image_row_pitch = 0;
+    imageDesc.image_slice_pitch = 0;
+    imageDesc.num_mip_levels = 0;
+    imageDesc.num_samples = 0;
+
+    auto retVal = CL_INVALID_VALUE;
+    cl_mem_flags flags = CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR;
+
+    auto surfaceFormat = Image::getSurfaceFormatFromTable(flags, &imageFormat, device->getHardwareInfo().capabilityTable.clVersionSupport);
+    auto image = std::unique_ptr<Image>(Image::create(
+        contextCl,
+        MemoryPropertiesParser::createMemoryProperties(flags, 0, 0),
+        flags,
+        0,
+        surfaceFormat,
+        &imageDesc,
+        imageDataDst,
+        retVal));
+    ASSERT_NE(nullptr, image.get());
+    EXPECT_FALSE(image->isMemObjZeroCopy());
+
+    auto bufferSrc = std::unique_ptr<Buffer>(Buffer::create(context,
+                                                            CL_MEM_READ_WRITE,
+                                                            imageSize,
+                                                            nullptr,
+                                                            retVal));
+    ASSERT_NE(nullptr, bufferSrc);
+
+    memcpy(image->getGraphicsAllocation()->getUnderlyingBuffer(), imageDataDst, imageSize);
+    memcpy(bufferSrc->getGraphicsAllocation()->getUnderlyingBuffer(), imageDataSrc, imageSize);
+
+    auto simulatedCsr = AUBFixture::getSimulatedCsr<FamilyType>();
+
+    simulatedCsr->writeMemory(*bufferSrc->getGraphicsAllocation());
+    simulatedCsr->writeMemory(*image->getGraphicsAllocation());
+
+    kernel->setArg(0, bufferSrc.get());
+    kernel->setArg(1, image.get());
+
+    int srcOffset = 0;
+    int dstOffset[4] = {0, 0, 0, 0};
+    int pitch[2] = {0, 0};
+
+    kernel->setArg(2, sizeof(srcOffset), &srcOffset);
+    kernel->setArg(3, sizeof(dstOffset), &dstOffset);
+    kernel->setArg(4, sizeof(pitch), &pitch);
+
+    retVal = this->pCmdQ->enqueueKernel(
+        kernel.get(),
+        workDim,
+        globalWorkOffset,
+        globalWorkSize,
+        localWorkSize,
+        numEventsInWaitList,
+        eventWaitList,
+        event);
+    EXPECT_EQ(CL_SUCCESS, retVal);
+
+    retVal = this->pCmdQ->finish();
+    EXPECT_EQ(CL_SUCCESS, retVal);
+
+    expectMemory<FamilyType>(reinterpret_cast<void *>(image->getGraphicsAllocation()->getGpuAddress()),
+                             imageDataSrc, imageSize);
+}

opencl/test/unit_test/fixtures/simple_arg_kernel_fixture.h

@@ -293,15 +293,23 @@ class BindlessKernelFixture : public ProgramFixture {
         // temporarily skip test in Debug
         GTEST_SKIP();
 #endif
-        cl_device_id deviceId = device;
-        cl_context clContext = context;
+        this->deviceCl = device;
+        this->contextCl = context;
+    }
+
+    void TearDown() override {
+        ProgramFixture::TearDown();
+    }
+
+    void createKernel(const std::string &programName, const std::string &kernelName) {
         DebugManager.flags.UseBindlessBuffers.set(true);
         DebugManager.flags.UseBindlessImages.set(true);
-
+        cl_device_id deviceId = deviceCl;
+        cl_context clContext = contextCl;
         CreateProgramFromBinary(
             clContext,
             &deviceId,
-            "bindless_copy_buffer");
+            programName);
         ASSERT_NE(nullptr, pProgram);
 
         retVal = pProgram->build(
@@ -315,19 +323,17 @@ class BindlessKernelFixture : public ProgramFixture {
 
         kernel.reset(Kernel::create<MockKernel>(
             pProgram,
-            *pProgram->getKernelInfo("StatefulCopyBuffer"),
+            *pProgram->getKernelInfo(kernelName.c_str()),
             &retVal));
         ASSERT_NE(nullptr, kernel);
         ASSERT_EQ(CL_SUCCESS, retVal);
     }
 
-    void TearDown() override {
-        ProgramFixture::TearDown();
-    }
-
     DebugManagerStateRestore restorer;
     std::unique_ptr<Kernel> kernel = nullptr;
     cl_int retVal = CL_SUCCESS;
+    ClDevice *deviceCl = nullptr;
+    Context *contextCl = nullptr;
 };
 
 } // namespace NEO

opencl/test/unit_test/kernel/kernel_arg_buffer_tests.cpp

@@ -268,3 +268,25 @@ TEST_F(KernelArgBufferTest, givenNotUsedBindlessBuffersAndBufferArgWhenPatchingS
     pKernel->patchBindlessSurfaceStateOffsets(sshOffset);
     EXPECT_EQ(0xdeadu, *patchLocation);
 }
+
+HWTEST_F(KernelArgBufferTest, givenUsedBindlessBuffersAndBuiltinKernelWhenPatchingSurfaceStateOffsetsThenOffsetIsNotPatched) {
+    using DataPortBindlessSurfaceExtendedMessageDescriptor = typename FamilyType::DataPortBindlessSurfaceExtendedMessageDescriptor;
+    DebugManagerStateRestore restorer;
+    DebugManager.flags.UseBindlessBuffers.set(1);
+
+    pKernelInfo->usesSsh = true;
+    pKernelInfo->requiresSshForBuffers = true;
+
+    auto crossThreadDataOffset = pKernelInfo->kernelArgInfo[0].kernelArgPatchInfoVector[0].crossthreadOffset;
+    pKernelInfo->kernelArgInfo[0].offsetHeap = 64;
+    pKernelInfo->kernelArgInfo[0].isBuffer = true;
+
+    auto patchLocation = reinterpret_cast<uint32_t *>(ptrOffset(pKernel->getCrossThreadData(), crossThreadDataOffset));
+    *patchLocation = 0xdead;
+
+    pKernel->isBuiltIn = true;
+
+    uint32_t sshOffset = 0x1000;
+    pKernel->patchBindlessSurfaceStateOffsets(sshOffset);
+    EXPECT_EQ(0xdeadu, *patchLocation);
+}

opencl/test/unit_test/test_files/copy_buffer_to_image.cl

@@ -0,0 +1,23 @@
+/*
+ * Copyright (C) 2020 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable
+
+__kernel void CopyBufferToImage3d(__global uchar *src,
+                                  __write_only image3d_t output,
+                                  int srcOffset,
+                                  int4 dstOffset,
+                                  uint2 Pitch) {
+    const uint x = get_global_id(0);
+    const uint y = get_global_id(1);
+    const uint z = get_global_id(2);
+
+    int4 dstCoord = (int4)(x, y, z, 0) + dstOffset;
+    uint LOffset = srcOffset + (y * Pitch.x) + (z * Pitch.y);
+
+    write_imageui(output, dstCoord, (uint4)(*(src + LOffset + x), 0, 0, 1));
+}