diff --git a/runtime/command_queue/enqueue_common.h b/runtime/command_queue/enqueue_common.h
index 5d074e368c..a2e833fab9 100644
--- a/runtime/command_queue/enqueue_common.h
+++ b/runtime/command_queue/enqueue_common.h
@@ -415,10 +415,6 @@ void CommandQueueHw<GfxFamily>::enqueueHandler(Surface **surfacesForResidency,
             waitUntilComplete(taskCount, flushStamp->peekStamp(), false);
         } else {
             waitUntilComplete(taskCount, flushStamp->peekStamp(), false);
-            for (auto sIt = surfacesForResidency, sE = surfacesForResidency + numSurfaceForResidency;
-                 sIt != sE; ++sIt) {
-                (*sIt)->setCompletionStamp(completionStamp, nullptr, nullptr);
-            }
             if (printfHandler) {
                 printfHandler->printEnqueueOutput();
             }
@@ -575,14 +571,6 @@ CompletionStamp CommandQueueHw<GfxFamily>::enqueueNonBlocked(
         dispatchFlags,
         *device);
 
-    for (auto surface : CreateRange(surfaces, surfaceCount)) {
-        surface->setCompletionStamp(completionStamp, device, this);
-    }
-
-    for (auto &dispatchInfo : multiDispatchInfo) {
-        dispatchInfo.getKernel()->updateWithCompletionStamp(commandStreamReceiver, &completionStamp);
-    }
-
     return completionStamp;
 }
 
diff --git a/runtime/device/device.h b/runtime/device/device.h
index 0afc7a41ad..928be71f31 100644
--- a/runtime/device/device.h
+++ b/runtime/device/device.h
@@ -131,6 +131,7 @@ class Device : public BaseObject<_cl_device_id> {
     bool getEnabled64kbPages();
     bool isSourceLevelDebuggerActive() const;
     SourceLevelDebugger *getSourceLevelDebugger() { return executionEnvironment->sourceLevelDebugger.get(); }
+    ExecutionEnvironment *getExecutionEnvironment() const { return executionEnvironment; }
 
   protected:
     Device() = delete;
diff --git a/runtime/helpers/task_information.cpp b/runtime/helpers/task_information.cpp
index 2bd72d0bba..6f371b149c 100644
--- a/runtime/helpers/task_information.cpp
+++ b/runtime/helpers/task_information.cpp
@@ -235,9 +235,6 @@ CompletionStamp &CommandComputeKernel::submit(uint32_t taskLevel, bool terminate
                                                       taskLevel,
                                                       dispatchFlags,
                                                       commandQueue.getDevice());
-    for (auto &surface : surfaces) {
-        surface->setCompletionStamp(completionStamp, nullptr, nullptr);
-    }
     commandQueue.waitUntilComplete(completionStamp.taskCount, completionStamp.flushStamp, false);
     if (printfHandler) {
         printfHandler.get()->printEnqueueOutput();
diff --git a/runtime/kernel/kernel.cpp b/runtime/kernel/kernel.cpp
index 5ba5609f73..59a54dcfc4 100644
--- a/runtime/kernel/kernel.cpp
+++ b/runtime/kernel/kernel.cpp
@@ -935,20 +935,6 @@ inline void Kernel::makeArgsResident(CommandStreamReceiver &commandStreamReceive
     }
 }
 
-void Kernel::updateWithCompletionStamp(CommandStreamReceiver &commandStreamReceiver, CompletionStamp *completionStamp) {
-    auto numArgs = kernelInfo.kernelArgInfo.size();
-    for (decltype(numArgs) argIndex = 0; argIndex < numArgs; argIndex++) {
-        if (kernelArguments[argIndex].object) {
-            if (Kernel::isMemObj(kernelArguments[argIndex].type)) {
-                auto clMem = const_cast<cl_mem>(static_cast<const _cl_mem *>(kernelArguments[argIndex].object));
-                auto memObj = castToObjectOrAbort<MemObj>(clMem);
-                DEBUG_BREAK_IF(memObj == nullptr);
-                memObj->setCompletionStamp(*completionStamp, nullptr, nullptr);
-            }
-        }
-    }
-}
-
 void Kernel::makeResident(CommandStreamReceiver &commandStreamReceiver) {
     if (privateSurface) {
         commandStreamReceiver.makeResident(*privateSurface);
diff --git a/runtime/kernel/kernel.h b/runtime/kernel/kernel.h
index 307d610a0a..e9ddf00503 100644
--- a/runtime/kernel/kernel.h
+++ b/runtime/kernel/kernel.h
@@ -292,7 +292,6 @@ class Kernel : public BaseObject<_cl_kernel> {
 
     //residency for kernel surfaces
     MOCKABLE_VIRTUAL void makeResident(CommandStreamReceiver &commandStreamReceiver);
-    void updateWithCompletionStamp(CommandStreamReceiver &commandStreamReceiver, CompletionStamp *completionStamp);
     MOCKABLE_VIRTUAL void getResidency(std::vector<Surface *> &dst);
     bool requiresCoherency();
     void resetSharedObjectsPatchAddresses();
diff --git a/runtime/mem_obj/buffer.cpp b/runtime/mem_obj/buffer.cpp
index 5a0bc20eb3..3fded6bada 100644
--- a/runtime/mem_obj/buffer.cpp
+++ b/runtime/mem_obj/buffer.cpp
@@ -438,7 +438,7 @@ Buffer *Buffer::createBufferHwFromDevice(const Device *device,
     DEBUG_BREAK_IF(nullptr == funcCreate);
     auto pBuffer = funcCreate(nullptr, flags, size, memoryStorage, hostPtr, gfxAllocation,
                               zeroCopy, isHostPtrSVM, isImageRedescribed);
-    pBuffer->device = const_cast<Device *>(device);
+    pBuffer->executionEnvironment = device->getExecutionEnvironment();
     return pBuffer;
 }
 
diff --git a/runtime/mem_obj/buffer.h b/runtime/mem_obj/buffer.h
index d6684faced..55048c8bf0 100644
--- a/runtime/mem_obj/buffer.h
+++ b/runtime/mem_obj/buffer.h
@@ -28,6 +28,7 @@
 
 namespace OCLRT {
 class Buffer;
+class Device;
 class MemoryManager;
 
 typedef Buffer *(*BufferCreatFunc)(Context *context,
diff --git a/runtime/mem_obj/buffer.inl b/runtime/mem_obj/buffer.inl
index d0f639345b..01f2e62e70 100644
--- a/runtime/mem_obj/buffer.inl
+++ b/runtime/mem_obj/buffer.inl
@@ -21,7 +21,7 @@
  */
 
 #include "hw_cmds.h"
-#include "runtime/device/device.h"
+#include "runtime/execution_environment/execution_environment.h"
 #include "runtime/helpers/surface_formats.h"
 #include "runtime/helpers/aligned_memory.h"
 #include "runtime/mem_obj/buffer.h"
@@ -46,7 +46,7 @@ void BufferHw<GfxFamily>::setArgStateful(void *memory) {
     using SURFACE_FORMAT = typename RENDER_SURFACE_STATE::SURFACE_FORMAT;
     using AUXILIARY_SURFACE_MODE = typename RENDER_SURFACE_STATE::AUXILIARY_SURFACE_MODE;
 
-    auto gmmHelper = device->getGmmHelper();
+    auto gmmHelper = executionEnvironment->getGmmHelper();
     auto surfaceState = reinterpret_cast<RENDER_SURFACE_STATE *>(memory);
     auto surfaceSize = alignUp(getSize(), 4);
 
diff --git a/runtime/mem_obj/image.inl b/runtime/mem_obj/image.inl
index 19d945e381..42f3d56878 100644
--- a/runtime/mem_obj/image.inl
+++ b/runtime/mem_obj/image.inl
@@ -21,7 +21,7 @@
  */
 
 #include "hw_cmds.h"
-#include "runtime/device/device.h"
+#include "runtime/execution_environment/execution_environment.h"
 #include "runtime/helpers/surface_formats.h"
 #include "runtime/helpers/aligned_memory.h"
 #include "runtime/mem_obj/image.h"
@@ -45,7 +45,7 @@ void ImageHw<GfxFamily>::setImageArg(void *memory, bool setAsMediaBlockImage, ui
     using SURFACE_FORMAT = typename RENDER_SURFACE_STATE::SURFACE_FORMAT;
     auto surfaceState = reinterpret_cast<RENDER_SURFACE_STATE *>(memory);
     auto gmm = getGraphicsAllocation()->gmm;
-    auto gmmHelper = device->getGmmHelper();
+    auto gmmHelper = executionEnvironment->getGmmHelper();
 
     auto imageCount = std::max(getImageDesc().image_depth, getImageDesc().image_array_size);
     if (imageCount == 0) {
@@ -220,7 +220,7 @@ void ImageHw<GfxFamily>::setMediaImageArg(void *memory) {
     using SURFACE_FORMAT = typename MEDIA_SURFACE_STATE::SURFACE_FORMAT;
     SURFACE_FORMAT surfaceFormat = MEDIA_SURFACE_STATE::SURFACE_FORMAT_Y8_UNORM_VA;
 
-    auto gmmHelper = device->getGmmHelper();
+    auto gmmHelper = executionEnvironment->getGmmHelper();
     auto surfaceState = reinterpret_cast<MEDIA_SURFACE_STATE *>(memory);
     *surfaceState = MEDIA_SURFACE_STATE::sInit();
 
diff --git a/runtime/mem_obj/mem_obj.cpp b/runtime/mem_obj/mem_obj.cpp
index 6baece5ad6..c18ad93c80 100644
--- a/runtime/mem_obj/mem_obj.cpp
+++ b/runtime/mem_obj/mem_obj.cpp
@@ -48,12 +48,11 @@ MemObj::MemObj(Context *context,
       memoryStorage(memoryStorage), hostPtr(hostPtr),
       isZeroCopy(zeroCopy), isHostPtrSVM(isHostPtrSVM), isObjectRedescribed(isObjectRedescribed),
       graphicsAllocation(gfxAllocation) {
-    completionStamp = {};
 
     if (context) {
         context->incRefInternal();
         memoryManager = context->getMemoryManager();
-        device = context->getDevice(0);
+        executionEnvironment = context->getDevice(0)->getExecutionEnvironment();
     }
 }
 
@@ -217,15 +216,6 @@ size_t MemObj::getSize() const {
     return size;
 }
 
-void MemObj::setCompletionStamp(CompletionStamp completionStamp, Device *pDevice, CommandQueue *pCmdQ) {
-    this->completionStamp = completionStamp;
-    cmdQueuePtr = pCmdQ;
-}
-
-CompletionStamp MemObj::getCompletionStamp() const {
-    return completionStamp;
-}
-
 void MemObj::setAllocatedMapPtr(void *allocatedMapPtr) {
     this->allocatedMapPtr = allocatedMapPtr;
 }
diff --git a/runtime/mem_obj/mem_obj.h b/runtime/mem_obj/mem_obj.h
index fdb2cf0ecc..2adddc6643 100644
--- a/runtime/mem_obj/mem_obj.h
+++ b/runtime/mem_obj/mem_obj.h
@@ -23,7 +23,6 @@
 #pragma once
 #include "runtime/api/cl_types.h"
 #include "runtime/helpers/base_object.h"
-#include "runtime/helpers/completion_stamp.h"
 #include "runtime/helpers/mipmap.h"
 #include "runtime/sharings/sharing.h"
 #include "runtime/mem_obj/map_operations_handler.h"
@@ -33,12 +32,11 @@
 #include <vector>
 
 namespace OCLRT {
+class ExecutionEnvironment;
 class GraphicsAllocation;
 struct KernelInfo;
 class MemoryManager;
 class Context;
-class CommandQueue;
-class Device;
 
 template <>
 struct OpenCLObjectMapper<_cl_mem> {
@@ -74,8 +72,6 @@ class MemObj : public BaseObject<_cl_mem> {
     bool getIsObjectRedescribed() const { return isObjectRedescribed; };
     size_t getSize() const;
     cl_mem_flags getFlags() const;
-    void setCompletionStamp(CompletionStamp completionStamp, Device *pDevice, CommandQueue *pCmdQ);
-    CompletionStamp getCompletionStamp() const;
 
     bool addMappedPtr(void *ptr, size_t ptrLength, cl_map_flags &mapFlags, MemObjSizeArray &size, MemObjOffsetArray &offset, uint32_t mipLevel);
     bool findMappedPtr(void *mappedPtr, MapInfo &outMapInfo) { return mapOperationsHandler.find(mappedPtr, outMapInfo); }
@@ -104,7 +100,6 @@ class MemObj : public BaseObject<_cl_mem> {
 
     virtual bool allowTiling() const { return false; }
 
-    CommandQueue *getAssociatedCommandQueue() { return cmdQueuePtr; }
     bool isImageFromImage() const { return isImageFromImageCreated; }
 
     void *getCpuAddressForMapping();
@@ -141,9 +136,7 @@ class MemObj : public BaseObject<_cl_mem> {
     size_t offset = 0;
     MemObj *associatedMemObject = nullptr;
     cl_uint refCount = 0;
-    CompletionStamp completionStamp;
-    CommandQueue *cmdQueuePtr = nullptr;
-    Device *device = nullptr;
+    ExecutionEnvironment *executionEnvironment = nullptr;
     bool isZeroCopy;
     bool isHostPtrSVM;
     bool isObjectRedescribed;
diff --git a/runtime/memory_manager/surface.h b/runtime/memory_manager/surface.h
index d53fe2220a..1d8d7b4af1 100644
--- a/runtime/memory_manager/surface.h
+++ b/runtime/memory_manager/surface.h
@@ -26,13 +26,12 @@
 #include "runtime/memory_manager/graphics_allocation.h"
 
 namespace OCLRT {
-
+class CommandQueue;
 class Surface {
   public:
     Surface(bool isCoherent = false) : IsCoherent(isCoherent) {}
     virtual ~Surface() = default;
     virtual void makeResident(CommandStreamReceiver &csr) = 0;
-    virtual void setCompletionStamp(CompletionStamp &cs, Device *pDevice, CommandQueue *pCmdQ) = 0;
     virtual Surface *duplicate() = 0;
     const bool IsCoherent;
 };
@@ -43,7 +42,6 @@ class NullSurface : public Surface {
     ~NullSurface() override = default;
 
     void makeResident(CommandStreamReceiver &csr) override{};
-    void setCompletionStamp(CompletionStamp &cs, Device *pDevice, CommandQueue *pCmdQ) override{};
     Surface *duplicate() override { return new NullSurface(); };
 };
 
@@ -67,10 +65,6 @@ class HostPtrSurface : public Surface {
         DEBUG_BREAK_IF(!gfxAllocation);
         csr.makeResidentHostPtrAllocation(gfxAllocation);
     }
-    void setCompletionStamp(CompletionStamp &cs, Device *pDevice, CommandQueue *pCmdQ) override {
-        DEBUG_BREAK_IF(!gfxAllocation);
-        gfxAllocation->taskCount = cs.taskCount;
-    }
     Surface *duplicate() override {
         return new HostPtrSurface(this->memoryPointer, this->surfaceSize, this->gfxAllocation);
     };
@@ -103,28 +97,25 @@ class HostPtrSurface : public Surface {
 
 class MemObjSurface : public Surface {
   public:
-    MemObjSurface(MemObj *memObj) : Surface(memObj->getGraphicsAllocation()->isCoherent()), memory_object(memObj) {
-        memory_object->retain();
+    MemObjSurface(MemObj *memObj) : Surface(memObj->getGraphicsAllocation()->isCoherent()), memObj(memObj) {
+        memObj->retain();
     }
     ~MemObjSurface() override {
-        memory_object->release();
-        memory_object = nullptr;
+        memObj->release();
+        memObj = nullptr;
     };
 
     void makeResident(CommandStreamReceiver &csr) override {
-        DEBUG_BREAK_IF(!memory_object);
-        csr.makeResident(*memory_object->getGraphicsAllocation());
-    }
-    void setCompletionStamp(CompletionStamp &cs, Device *pDevice, CommandQueue *pCmdQ) override {
-        DEBUG_BREAK_IF(!memory_object);
-        memory_object->setCompletionStamp(cs, pDevice, pCmdQ);
+        DEBUG_BREAK_IF(!memObj);
+        csr.makeResident(*memObj->getGraphicsAllocation());
     }
+
     Surface *duplicate() override {
-        return new MemObjSurface(this->memory_object);
+        return new MemObjSurface(this->memObj);
     };
 
   protected:
-    class MemObj *memory_object;
+    class MemObj *memObj;
 };
 
 class GeneralSurface : public Surface {
@@ -137,9 +128,6 @@ class GeneralSurface : public Surface {
     void makeResident(CommandStreamReceiver &csr) override {
         csr.makeResident(*gfxAllocation);
     };
-    void setCompletionStamp(CompletionStamp &cs, Device *pDevice, CommandQueue *pCmdQ) override {
-        gfxAllocation->taskCount = cs.taskCount;
-    };
     Surface *duplicate() override { return new GeneralSurface(gfxAllocation); };
 
   protected:
diff --git a/unit_tests/command_queue/enqueue_fill_buffer_tests.cpp b/unit_tests/command_queue/enqueue_fill_buffer_tests.cpp
index dd0b6a4f11..6e9c32a798 100644
--- a/unit_tests/command_queue/enqueue_fill_buffer_tests.cpp
+++ b/unit_tests/command_queue/enqueue_fill_buffer_tests.cpp
@@ -56,15 +56,6 @@ HWTEST_F(EnqueueFillBufferCmdTests, bumpsTaskLevel) {
     EXPECT_GT(pCmdQ->taskLevel, taskLevelBefore);
 }
 
-HWTEST_F(EnqueueFillBufferCmdTests, setsBufferCompletionStamp) {
-    EnqueueFillBufferHelper<>::enqueueFillBuffer(pCmdQ, buffer);
-    auto deviceEngineType = pDevice->getEngineType();
-    auto &commandStreamReceiver = pDevice->getCommandStreamReceiver();
-    EXPECT_EQ(commandStreamReceiver.peekTaskCount(), buffer->getCompletionStamp().taskCount);
-    EXPECT_EQ(0u, buffer->getCompletionStamp().deviceOrdinal);
-    EXPECT_EQ(deviceEngineType, buffer->getCompletionStamp().engineType);
-}
-
 HWTEST_F(EnqueueFillBufferCmdTests, addsCommands) {
     auto usedCmdBufferBefore = pCS->getUsed();
 
diff --git a/unit_tests/command_queue/enqueue_kernel_tests.cpp b/unit_tests/command_queue/enqueue_kernel_tests.cpp
index a243f9619a..e86b7045bf 100644
--- a/unit_tests/command_queue/enqueue_kernel_tests.cpp
+++ b/unit_tests/command_queue/enqueue_kernel_tests.cpp
@@ -58,40 +58,6 @@ TEST_F(EnqueueKernelTest, clEnqueueNDRangeKernel_null_kernel) {
     EXPECT_EQ(CL_INVALID_KERNEL, retVal);
 }
 
-TEST_F(EnqueueKernelTest, clEnqueueNDRangeKernel_SetCompletionStampTaskCountOnBuffersAndDcFlushOnMapBufferIfItsTaskCountIsGreaterThanLatestDcFlushTaskCount) {
-    const size_t n = 512;
-    size_t globalWorkSize[3] = {n, 1, 1};
-    size_t localWorkSize[3] = {256, 1, 1};
-    cl_int retVal = CL_SUCCESS;
-    CommandQueue *pCmdQ2 = createCommandQueue(pDevice, 0);
-
-    auto b0 = clCreateBuffer(context, 0, n * sizeof(float), nullptr, nullptr);
-    auto b1 = clCreateBuffer(context, 0, n * sizeof(float), nullptr, nullptr);
-
-    retVal = clSetKernelArg(pKernel, 0, sizeof(cl_mem), &b0);
-    EXPECT_EQ(CL_SUCCESS, retVal);
-    retVal = clSetKernelArg(pKernel, 1, sizeof(cl_mem), &b1);
-    EXPECT_EQ(CL_SUCCESS, retVal);
-    retVal = clEnqueueNDRangeKernel(pCmdQ2, pKernel, 1, nullptr, globalWorkSize, localWorkSize, 0, nullptr, nullptr);
-    EXPECT_EQ(CL_SUCCESS, retVal);
-    auto taskCount0 = castToObject<MemObj>(b0)->getCompletionStamp().taskCount;
-    auto taskCount1 = castToObject<MemObj>(b1)->getCompletionStamp().taskCount;
-
-    EXPECT_EQ(1u, taskCount0);
-    EXPECT_EQ(1u, taskCount1);
-
-    auto ptrResult = clEnqueueMapBuffer(pCmdQ, b1, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, 8, 0, nullptr, nullptr, &retVal);
-    EXPECT_EQ(CL_SUCCESS, retVal);
-    clEnqueueUnmapMemObject(pCmdQ, b1, ptrResult, 0, nullptr, nullptr);
-
-    retVal = clReleaseMemObject(b0);
-    EXPECT_EQ(CL_SUCCESS, retVal);
-    retVal = clReleaseMemObject(b1);
-    EXPECT_EQ(CL_SUCCESS, retVal);
-    clReleaseCommandQueue(pCmdQ2);
-    EXPECT_EQ(CL_SUCCESS, retVal);
-}
-
 TEST_F(EnqueueKernelTest, givenKernelWhenAllArgsAreSetThenClEnqueueNDRangeKernelReturnsSuccess) {
     const size_t n = 512;
     size_t globalWorkSize[3] = {n, 1, 1};
diff --git a/unit_tests/command_queue/enqueue_read_buffer_tests.cpp b/unit_tests/command_queue/enqueue_read_buffer_tests.cpp
index d3b4f18aad..136c7b1169 100644
--- a/unit_tests/command_queue/enqueue_read_buffer_tests.cpp
+++ b/unit_tests/command_queue/enqueue_read_buffer_tests.cpp
@@ -85,9 +85,6 @@ HWCMDTEST_F(IGFX_GEN8_CORE, EnqueueReadBufferTypeTest, GPGPUWalker) {
     EXPECT_NE(0u, cmd->getIndirectDataLength());
     EXPECT_FALSE(cmd->getIndirectParameterEnable());
 
-    // Verify srcBuffer internal state
-    EXPECT_EQ(nullptr, srcBuffer->getAssociatedCommandQueue());
-
     // Compute the SIMD lane mask
     size_t simd =
         cmd->getSimdSize() == GPGPU_WALKER::SIMD_SIZE_SIMD32 ? 32 : cmd->getSimdSize() == GPGPU_WALKER::SIMD_SIZE_SIMD16 ? 16 : 8;
diff --git a/unit_tests/event/event_tests.cpp b/unit_tests/event/event_tests.cpp
index 78eb299ca8..67fc8bd67b 100644
--- a/unit_tests/event/event_tests.cpp
+++ b/unit_tests/event/event_tests.cpp
@@ -413,7 +413,7 @@ TEST_F(EventTest, givenEventContainingCommandQueueWhenItsStatusIsUpdatedToComple
 class SurfaceMock : public Surface {
   public:
     SurfaceMock() {
-        resident = nonResident = completionStamp = 0;
+        resident = nonResident = 0;
     };
     ~SurfaceMock() override{};
 
@@ -427,13 +427,6 @@ class SurfaceMock : public Surface {
             csr.makeResident(*graphicsAllocation);
         }
     };
-    void setCompletionStamp(CompletionStamp &cs, Device *pDevice, CommandQueue *pCmdQ) override {
-        if (parent) {
-            parent->completionStamp++;
-        } else {
-            completionStamp++;
-        }
-    };
     Surface *duplicate() override {
         return new SurfaceMock(this);
     };
@@ -441,7 +434,6 @@ class SurfaceMock : public Surface {
     SurfaceMock *parent = nullptr;
     std::atomic<uint32_t> resident;
     std::atomic<uint32_t> nonResident;
-    std::atomic<uint32_t> completionStamp;
 
     GraphicsAllocation *graphicsAllocation = nullptr;
 
@@ -486,7 +478,6 @@ TEST_F(InternalsEventTest, processBlockedCommandsKernelOperation) {
     EXPECT_EQ(surface->resident, 1u);
     EXPECT_FALSE(surface->graphicsAllocation->isResident());
     delete surface->graphicsAllocation;
-    EXPECT_EQ(surface->completionStamp, 1u);
 }
 
 TEST_F(InternalsEventTest, processBlockedCommandsAbortKernelOperation) {
diff --git a/unit_tests/gtpin/gtpin_tests.cpp b/unit_tests/gtpin/gtpin_tests.cpp
index 04585acbd6..616cb36a05 100644
--- a/unit_tests/gtpin/gtpin_tests.cpp
+++ b/unit_tests/gtpin/gtpin_tests.cpp
@@ -788,16 +788,6 @@ TEST_F(GTPinTests, givenInitializedGTPinInterfaceWhenKernelIsExecutedThenGTPinCa
     EXPECT_EQ(prevCount14 + 2, CommandBufferCompleteCallbackCount);
     EXPECT_EQ(prevCount24 + 2, CommandBufferCompleteCallbackCount);
 
-    auto taskCount0 = castToObject<MemObj>(buff10)->getCompletionStamp().taskCount;
-    auto taskCount1 = castToObject<MemObj>(buff11)->getCompletionStamp().taskCount;
-    EXPECT_EQ(1u, taskCount0);
-    EXPECT_EQ(1u, taskCount1);
-
-    taskCount0 = castToObject<MemObj>(buff20)->getCompletionStamp().taskCount;
-    taskCount1 = castToObject<MemObj>(buff21)->getCompletionStamp().taskCount;
-    EXPECT_EQ(2u, taskCount0);
-    EXPECT_EQ(2u, taskCount1);
-
     // Cleanup
     retVal = clReleaseKernel(kernel1);
     EXPECT_EQ(CL_SUCCESS, retVal);
diff --git a/unit_tests/kernel/kernel_tests.cpp b/unit_tests/kernel/kernel_tests.cpp
index 5936fffdf6..d4377d60cf 100644
--- a/unit_tests/kernel/kernel_tests.cpp
+++ b/unit_tests/kernel/kernel_tests.cpp
@@ -508,7 +508,6 @@ TEST_F(KernelPrivateSurfaceTest, testPrivateSurface) {
     EXPECT_EQ(1u, csr->residency.size());
 
     csr->makeSurfacePackNonResident(nullptr, false);
-    pKernel->updateWithCompletionStamp(*csr.get(), nullptr);
     EXPECT_EQ(0u, csr->residency.size());
 
     delete pKernel;
diff --git a/unit_tests/mem_obj/buffer_tests.cpp b/unit_tests/mem_obj/buffer_tests.cpp
index b5c67b8629..255ab1b74a 100644
--- a/unit_tests/mem_obj/buffer_tests.cpp
+++ b/unit_tests/mem_obj/buffer_tests.cpp
@@ -547,37 +547,6 @@ TEST_P(NoHostPtr, GivenNoHostPtrWhenHwBufferCreationFailsThenReturnNullptr) {
     }
 }
 
-TEST_P(NoHostPtr, completionStamp) {
-    auto buffer = Buffer::create(
-        context.get(),
-        flags,
-        g_scTestBufferSizeInBytes,
-        nullptr,
-        retVal);
-
-    FlushStamp expectedFlushstamp = 0;
-
-    ASSERT_EQ(CL_SUCCESS, retVal);
-    ASSERT_NE(nullptr, buffer);
-    EXPECT_EQ(0u, buffer->getCompletionStamp().taskCount);
-    EXPECT_EQ(expectedFlushstamp, buffer->getCompletionStamp().flushStamp);
-    EXPECT_EQ(0u, buffer->getCompletionStamp().deviceOrdinal);
-    EXPECT_EQ(0u, buffer->getCompletionStamp().engineType);
-
-    CompletionStamp completionStamp;
-    completionStamp.taskCount = 42;
-    completionStamp.deviceOrdinal = 43;
-    completionStamp.engineType = EngineType::ENGINE_RCS;
-    completionStamp.flushStamp = 5;
-    buffer->setCompletionStamp(completionStamp, nullptr, nullptr);
-    EXPECT_EQ(completionStamp.taskCount, buffer->getCompletionStamp().taskCount);
-    EXPECT_EQ(completionStamp.flushStamp, buffer->getCompletionStamp().flushStamp);
-    EXPECT_EQ(completionStamp.deviceOrdinal, buffer->getCompletionStamp().deviceOrdinal);
-    EXPECT_EQ(completionStamp.engineType, buffer->getCompletionStamp().engineType);
-
-    delete buffer;
-}
-
 TEST_P(NoHostPtr, WithUseHostPtr_returnsError) {
     auto buffer = Buffer::create(
         context.get(),
diff --git a/unit_tests/mem_obj/image_tests.cpp b/unit_tests/mem_obj/image_tests.cpp
index 094020213e..231a61a343 100644
--- a/unit_tests/mem_obj/image_tests.cpp
+++ b/unit_tests/mem_obj/image_tests.cpp
@@ -447,33 +447,6 @@ TEST_P(CreateImageNoHostPtr, getImageDesc) {
     delete image;
 }
 
-TEST_P(CreateImageNoHostPtr, completionStamp) {
-    auto image = createImageWithFlags(flags);
-
-    FlushStamp expectedFlushStamp = 0;
-
-    ASSERT_EQ(CL_SUCCESS, retVal);
-    ASSERT_NE(nullptr, image);
-    ASSERT_EQ(true, image->isMemObjZeroCopy());
-    EXPECT_EQ(0u, image->getCompletionStamp().taskCount);
-    EXPECT_EQ(expectedFlushStamp, image->getCompletionStamp().flushStamp);
-    EXPECT_EQ(0u, image->getCompletionStamp().deviceOrdinal);
-    EXPECT_EQ(0u, image->getCompletionStamp().engineType);
-
-    CompletionStamp completionStamp;
-    completionStamp.taskCount = 42;
-    completionStamp.deviceOrdinal = 43;
-    completionStamp.engineType = EngineType::ENGINE_RCS;
-    completionStamp.flushStamp = 5;
-    image->setCompletionStamp(completionStamp, nullptr, nullptr);
-    EXPECT_EQ(completionStamp.taskCount, image->getCompletionStamp().taskCount);
-    EXPECT_EQ(completionStamp.flushStamp, image->getCompletionStamp().flushStamp);
-    EXPECT_EQ(completionStamp.deviceOrdinal, image->getCompletionStamp().deviceOrdinal);
-    EXPECT_EQ(completionStamp.engineType, image->getCompletionStamp().engineType);
-
-    delete image;
-}
-
 TEST_P(CreateImageNoHostPtr, withImageGraphicsAllocationReportsImageType) {
     auto image = createImageWithFlags(flags);
 
diff --git a/unit_tests/program/program_tests.cpp b/unit_tests/program/program_tests.cpp
index 6bb86b58d4..55b75a3af9 100644
--- a/unit_tests/program/program_tests.cpp
+++ b/unit_tests/program/program_tests.cpp
@@ -1416,7 +1416,6 @@ HWTEST_F(PatchTokenTests, givenKernelRequiringConstantAllocationWhenMakeResident
 
     EXPECT_EQ(*pDst, reinterpret_cast<uintptr_t>(constBuffGpuAddr));
 
-    pKernel->updateWithCompletionStamp(*pCommandStreamReceiver, nullptr);
     pCommandStreamReceiver->makeSurfacePackNonResident(nullptr, false);
     EXPECT_EQ(0u, pCommandStreamReceiver->residency.size());