Add new simple kernel and method to adjust WALKER command parameters

Change-Id: Id0591908353ca744c44c7bab4e27db8332289a68

runtime/command_queue/gpgpu_walker.h

@@ -201,6 +201,11 @@ class GpgpuWalkerHelper {
 
     static void dispatchOnDeviceWaitlistSemaphores(LinearStream *commandStream, Device &currentDevice,
                                                    cl_uint numEventsInWaitList, const cl_event *eventWaitList);
+
+    static void adjustWalkerData(LinearStream *commandStream,
+                                 WALKER_TYPE<GfxFamily> *walkerCmd,
+                                 const Kernel &kernel,
+                                 const DispatchInfo &dispatchInfo);
 };
 
 template <typename GfxFamily>

runtime/command_queue/gpgpu_walker_base.inl

@@ -189,4 +189,11 @@ void GpgpuWalkerHelper<GfxFamily>::setupTimestampPacket(
     }
 }
 
+template <typename GfxFamily>
+void GpgpuWalkerHelper<GfxFamily>::adjustWalkerData(LinearStream *commandStream,
+                                                    WALKER_TYPE<GfxFamily> *walkerCmd,
+                                                    const Kernel &kernel,
+                                                    const DispatchInfo &dispatchInfo) {
+}
+
 } // namespace OCLRT

runtime/command_queue/hardware_interface.inl

@@ -226,6 +226,8 @@ void HardwareInterface<GfxFamily>::dispatchWalker(
                                                                localIdsGenerationByRuntime, inlineDataProgrammingRequired,
                                                                *kernel.getKernelInfo().patchInfo.threadPayload);
 
+        GpgpuWalkerHelper<GfxFamily>::adjustWalkerData(commandStream, walkerCmd, kernel, dispatchInfo);
+
         dispatchWorkarounds(commandStream, commandQueue, kernel, false);
         currentDispatchIndex++;
     }

unit_tests/fixtures/simple_arg_kernel_fixture.h

@@ -9,6 +9,7 @@
 #include "gtest/gtest.h"
 #include "CL/cl.h"
 #include "runtime/device/device.h"
+#include "runtime/helpers/array_count.h"
 #include "runtime/helpers/file_io.h"
 #include "runtime/kernel/kernel.h"
 #include "runtime/program/program.h"
@@ -187,7 +188,7 @@ class SimpleKernelFixture : public ProgramFixture {
   public:
     using ProgramFixture::SetUp;
     SimpleKernelFixture() {
-        kernelsCount = sizeof(kernels) / sizeof(Kernel *);
+        kernelsCount = arrayCount(kernels);
     }
 
   protected:
@@ -212,14 +213,14 @@ class SimpleKernelFixture : public ProgramFixture {
             false);
         ASSERT_EQ(CL_SUCCESS, retVal);
 
-        for (uint32_t i = 0; i < kernelsCount; i++) {
+        for (size_t i = 0; i < kernelsCount; i++) {
             if ((1 << i) & kernelIds) {
                 std::string kernelName("simple_kernel_");
                 kernelName.append(std::to_string(i));
-                kernels[i] = Kernel::create<MockKernel>(
+                kernels[i].reset(Kernel::create<MockKernel>(
                     pProgram,
                     *pProgram->getKernelInfo(kernelName.c_str()),
-                    &retVal);
+                    &retVal));
                 ASSERT_NE(nullptr, kernels[i]);
                 ASSERT_EQ(CL_SUCCESS, retVal);
             }
@@ -227,19 +228,18 @@ class SimpleKernelFixture : public ProgramFixture {
     }
 
     virtual void TearDown() {
-        for (uint32_t i = 0; i < kernelsCount; i++) {
+        for (size_t i = 0; i < kernelsCount; i++) {
             if (kernels[i]) {
-                delete kernels[i];
+                kernels[i].reset(nullptr);
-                kernels[i] = nullptr;
             }
         }
 
         ProgramFixture::TearDown();
     }
 
-    uint32_t kernelsCount;
+    size_t kernelsCount;
     cl_int retVal = CL_SUCCESS;
-    Kernel *kernels[5] = {};
+    std::unique_ptr<Kernel> kernels[6] = {};
     uint32_t kernelIds = 0;
 };
 

unit_tests/kernel/kernel_tests.cpp

@@ -395,6 +395,25 @@ TEST_F(KernelFromBinaryTests, BuiltInIsSetToFalseForRegularKernels) {
     EXPECT_FALSE(isBuiltIn);
 
     delete pKernel;
+    pKernel = nullptr;
+
+    pKernelInfo = pProgram->getKernelInfo("simple_kernel_5");
+
+    pKernel = Kernel::create(
+        pProgram,
+        *pKernelInfo,
+        &retVal);
+
+    ASSERT_EQ(CL_SUCCESS, retVal);
+    ASSERT_NE(nullptr, pKernel);
+
+    // get builtIn property
+    isBuiltIn = pKernel->isBuiltIn;
+
+    EXPECT_FALSE(isBuiltIn);
+
+    delete pKernel;
+    pKernel = nullptr;
 }
 
 TEST(PatchInfo, Constructor) {

unit_tests/test_files/simple_kernels.cl

@@ -5,6 +5,8 @@
  *
  */
 
+#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
+
 __kernel void simple_kernel_0(
     const uint arg0,
     const float arg1,
@@ -42,3 +44,7 @@ __kernel void simple_kernel_3(
 
 __kernel void simple_kernel_4() {
 }
+
+__kernel void simple_kernel_5(__global uint *dst) {
+    atomic_inc(dst);
+}