From c9e25a0fafd4849a2720d9a1a867933b2f814e60 Mon Sep 17 00:00:00 2001
From: "Mrozek, Michal" <michal.mrozek@intel.com>
Date: Wed, 31 Jul 2019 11:27:24 +0200
Subject: [PATCH] Mark constant arguments as read only in kernel argument info.

Change-Id: Icf49d9da060a144bc73e580ab77245e580e4812d
Signed-off-by: Mrozek, Michal <michal.mrozek@intel.com>
---
 runtime/program/kernel_arg_info.h       |  1 +
 runtime/program/kernel_info.cpp         |  1 +
 unit_tests/kernel/kernel_tests.cpp      | 20 ++++++++++++++++++++
 unit_tests/test_files/simple_kernels.cl |  4 ++--
 4 files changed, 24 insertions(+), 2 deletions(-)

diff --git a/runtime/program/kernel_arg_info.h b/runtime/program/kernel_arg_info.h
index 59ece15791..e42b3d917a 100644
--- a/runtime/program/kernel_arg_info.h
+++ b/runtime/program/kernel_arg_info.h
@@ -36,6 +36,7 @@ struct KernelArgInfo {
     bool isDeviceQueue = false;
     bool isBuffer = false;
     bool pureStatefulBufferAccess = false;
+    bool isReadOnly = false;
     uint32_t samplerArgumentType = 0;
     uint32_t offsetImgWidth = undefinedOffset;
     uint32_t offsetImgHeight = undefinedOffset;
diff --git a/runtime/program/kernel_info.cpp b/runtime/program/kernel_info.cpp
index c9b84de29a..8636e484ca 100644
--- a/runtime/program/kernel_info.cpp
+++ b/runtime/program/kernel_info.cpp
@@ -328,6 +328,7 @@ void KernelInfo::storeKernelArgument(
     usesSsh |= true;
     storeKernelArgPatchInfo(argNum, pStatelessConstMemObjKernelArg->DataParamSize, pStatelessConstMemObjKernelArg->DataParamOffset, 0, offsetSSH);
     kernelArgInfo[argNum].isBuffer = true;
+    kernelArgInfo[argNum].isReadOnly = true;
     patchInfo.statelessGlobalMemObjKernelArgs.push_back(reinterpret_cast<const SPatchStatelessGlobalMemoryObjectKernelArgument *>(pStatelessConstMemObjKernelArg));
 }
 
diff --git a/unit_tests/kernel/kernel_tests.cpp b/unit_tests/kernel/kernel_tests.cpp
index 03d323626e..0c583f8837 100644
--- a/unit_tests/kernel/kernel_tests.cpp
+++ b/unit_tests/kernel/kernel_tests.cpp
@@ -391,6 +391,26 @@ TEST_F(KernelFromBinaryTests, BuiltInIsSetToFalseForRegularKernels) {
     delete pKernel;
 }
 
+TEST_F(KernelFromBinaryTests, givenArgumentDeclaredAsConstantWhenKernelIsCreatedThenArgumentIsMarkedAsReadOnly) {
+    cl_device_id device = pDevice;
+
+    CreateProgramFromBinary<Program>(pContext, &device, "simple_kernels");
+
+    ASSERT_NE(nullptr, pProgram);
+    retVal = pProgram->build(
+        1,
+        &device,
+        nullptr,
+        nullptr,
+        nullptr,
+        false);
+
+    ASSERT_EQ(CL_SUCCESS, retVal);
+
+    auto pKernelInfo = pProgram->getKernelInfo("simple_kernel_6");
+    EXPECT_TRUE(pKernelInfo->kernelArgInfo[1].isReadOnly);
+}
+
 TEST(PatchInfo, Constructor) {
     PatchInfo patchInfo;
     EXPECT_EQ(nullptr, patchInfo.interfaceDescriptorDataLoad);
diff --git a/unit_tests/test_files/simple_kernels.cl b/unit_tests/test_files/simple_kernels.cl
index 9944c4b2df..536b1e9c47 100644
--- a/unit_tests/test_files/simple_kernels.cl
+++ b/unit_tests/test_files/simple_kernels.cl
@@ -19,7 +19,7 @@ __kernel void simple_kernel_0(
 }
 
 __kernel void simple_kernel_1(
-    __global uint *src,
+    __global const uint *src,
     const uint arg1,
     __global uint *dst) {
 
@@ -63,7 +63,7 @@ __kernel void simple_kernel_5(__global uint *dst) {
 }
 
 #define SIMPLE_KERNEL_6_ARRAY_SIZE 32
-__kernel void simple_kernel_6(__global uint *dst, __global uint2 *src, uint scalar, uint maxIterations, uint maxIterations2) {
+__kernel void simple_kernel_6(__global uint *dst, __constant uint2 *src, uint scalar, uint maxIterations, uint maxIterations2) {
     __private uint2 array[SIMPLE_KERNEL_6_ARRAY_SIZE];
     __private uint2 sum;
     __private size_t gid = get_global_id(0);