diff --git a/opencl/source/kernel/multi_device_kernel.cpp b/opencl/source/kernel/multi_device_kernel.cpp
index e104727f85..350c863765 100644
--- a/opencl/source/kernel/multi_device_kernel.cpp
+++ b/opencl/source/kernel/multi_device_kernel.cpp
@@ -9,11 +9,22 @@
 namespace NEO {
 
 MultiDeviceKernel::~MultiDeviceKernel() {
-    kernel->decRefInternal();
+    for (auto &pKernel : kernels) {
+        if (pKernel) {
+            pKernel->decRefInternal();
+        }
+    }
 }
-MultiDeviceKernel::MultiDeviceKernel(Kernel *pKernel) : kernel(pKernel) {
-    pKernel->incRefInternal();
-    pKernel->setMultiDeviceKernel(this);
+MultiDeviceKernel::MultiDeviceKernel(KernelVectorType kernelVector) : kernels(std::move(kernelVector)) {
+    for (auto &pKernel : kernels) {
+        if (pKernel) {
+            if (!defaultKernel) {
+                defaultKernel = kernels[(*pKernel->getDevices().begin())->getRootDeviceIndex()];
+            }
+            pKernel->incRefInternal();
+            pKernel->setMultiDeviceKernel(this);
+        }
+    }
 };
 
 } // namespace NEO
diff --git a/opencl/source/kernel/multi_device_kernel.h b/opencl/source/kernel/multi_device_kernel.h
index d0390a0899..126a186598 100644
--- a/opencl/source/kernel/multi_device_kernel.h
+++ b/opencl/source/kernel/multi_device_kernel.h
@@ -14,48 +14,59 @@ struct OpenCLObjectMapper<_cl_kernel> {
     typedef class MultiDeviceKernel DerivedType;
 };
 
+using KernelVectorType = StackVec<Kernel *, 4>;
+
 class MultiDeviceKernel : public BaseObject<_cl_kernel> {
   public:
     static const cl_ulong objectMagic = 0x3284ADC8EA0AFE25LL;
 
     ~MultiDeviceKernel() override;
-    MultiDeviceKernel(Kernel *pKernel);
+    MultiDeviceKernel(KernelVectorType kernelVector);
 
-    Kernel *getKernel(uint32_t rootDeviceIndex) const { return kernel; }
-    Kernel *getDefaultKernel() const { return kernel; }
+    Kernel *getKernel(uint32_t rootDeviceIndex) const { return kernels[rootDeviceIndex]; }
+    Kernel *getDefaultKernel() const { return defaultKernel; }
 
     template <typename kernel_t = Kernel, typename program_t = Program, typename multi_device_kernel_t = MultiDeviceKernel>
     static multi_device_kernel_t *create(program_t *program, const KernelInfoContainer &kernelInfos, cl_int *errcodeRet) {
+        KernelVectorType kernels{};
+        kernels.resize(program->getMaxRootDeviceIndex() + 1);
 
-        auto pKernel = Kernel::create<kernel_t, program_t>(program, kernelInfos, errcodeRet);
-        auto pMultiDeviceKernel = new multi_device_kernel_t(pKernel);
+        for (auto &pDevice : program->getDevices()) {
+            auto rootDeviceIndex = pDevice->getRootDeviceIndex();
+            if (kernels[rootDeviceIndex]) {
+                continue;
+            }
+            kernels[rootDeviceIndex] = Kernel::create<kernel_t, program_t>(program, kernelInfos, errcodeRet);
+        }
+        auto pMultiDeviceKernel = new multi_device_kernel_t(std::move(kernels));
 
         return pMultiDeviceKernel;
     }
 
-    cl_int cloneKernel(Kernel *pSourceKernel) { return kernel->cloneKernel(pSourceKernel); }
-    const std::vector<Kernel::SimpleKernelArgInfo> &getKernelArguments() const { return kernel->getKernelArguments(); }
-    cl_int checkCorrectImageAccessQualifier(cl_uint argIndex, size_t argSize, const void *argValue) const { return kernel->checkCorrectImageAccessQualifier(argIndex, argSize, argValue); }
-    void unsetArg(uint32_t argIndex) { return kernel->unsetArg(argIndex); }
-    cl_int setArg(uint32_t argIndex, size_t argSize, const void *argVal) { return kernel->setArg(argIndex, argSize, argVal); }
-    cl_int getInfo(cl_kernel_info paramName, size_t paramValueSize, void *paramValue, size_t *paramValueSizeRet) const { return kernel->getInfo(paramName, paramValueSize, paramValue, paramValueSizeRet); }
-    cl_int getArgInfo(cl_uint argIndx, cl_kernel_arg_info paramName, size_t paramValueSize, void *paramValue, size_t *paramValueSizeRet) const { return kernel->getArgInfo(argIndx, paramName, paramValueSize, paramValue, paramValueSizeRet); }
-    const ClDeviceVector &getDevices() const { return kernel->getDevices(); }
-    size_t getKernelArgsNumber() const { return kernel->getKernelArgsNumber(); }
-    Context &getContext() const { return kernel->getContext(); }
-    cl_int setArgSvmAlloc(uint32_t argIndex, void *svmPtr, GraphicsAllocation *svmAlloc) { return kernel->setArgSvmAlloc(argIndex, svmPtr, svmAlloc); }
-    bool getHasIndirectAccess() const { return kernel->getHasIndirectAccess(); }
-    void setUnifiedMemoryProperty(cl_kernel_exec_info infoType, bool infoValue) { return kernel->setUnifiedMemoryProperty(infoType, infoValue); }
-    void setSvmKernelExecInfo(GraphicsAllocation *argValue) { return kernel->setSvmKernelExecInfo(argValue); }
-    void clearSvmKernelExecInfo() { return kernel->clearSvmKernelExecInfo(); }
-    void setUnifiedMemoryExecInfo(GraphicsAllocation *argValue) { return kernel->setUnifiedMemoryExecInfo(argValue); }
-    void clearUnifiedMemoryExecInfo() { return kernel->clearUnifiedMemoryExecInfo(); }
-    int setKernelThreadArbitrationPolicy(uint32_t propertyValue) { return kernel->setKernelThreadArbitrationPolicy(propertyValue); }
-    cl_int setKernelExecutionType(cl_execution_info_kernel_type_intel executionType) { return kernel->setKernelExecutionType(executionType); }
-    int32_t setAdditionalKernelExecInfoWithParam(uint32_t paramName, size_t paramValueSize, const void *paramValue) { return kernel->setAdditionalKernelExecInfoWithParam(paramName, paramValueSize, paramValue); }
+    cl_int cloneKernel(Kernel *pSourceKernel) { return defaultKernel->cloneKernel(pSourceKernel); }
+    const std::vector<Kernel::SimpleKernelArgInfo> &getKernelArguments() const { return defaultKernel->getKernelArguments(); }
+    cl_int checkCorrectImageAccessQualifier(cl_uint argIndex, size_t argSize, const void *argValue) const { return defaultKernel->checkCorrectImageAccessQualifier(argIndex, argSize, argValue); }
+    void unsetArg(uint32_t argIndex) { return defaultKernel->unsetArg(argIndex); }
+    cl_int setArg(uint32_t argIndex, size_t argSize, const void *argVal) { return defaultKernel->setArg(argIndex, argSize, argVal); }
+    cl_int getInfo(cl_kernel_info paramName, size_t paramValueSize, void *paramValue, size_t *paramValueSizeRet) const { return defaultKernel->getInfo(paramName, paramValueSize, paramValue, paramValueSizeRet); }
+    cl_int getArgInfo(cl_uint argIndx, cl_kernel_arg_info paramName, size_t paramValueSize, void *paramValue, size_t *paramValueSizeRet) const { return defaultKernel->getArgInfo(argIndx, paramName, paramValueSize, paramValue, paramValueSizeRet); }
+    const ClDeviceVector &getDevices() const { return defaultKernel->getDevices(); }
+    size_t getKernelArgsNumber() const { return defaultKernel->getKernelArgsNumber(); }
+    Context &getContext() const { return defaultKernel->getContext(); }
+    cl_int setArgSvmAlloc(uint32_t argIndex, void *svmPtr, GraphicsAllocation *svmAlloc) { return defaultKernel->setArgSvmAlloc(argIndex, svmPtr, svmAlloc); }
+    bool getHasIndirectAccess() const { return defaultKernel->getHasIndirectAccess(); }
+    void setUnifiedMemoryProperty(cl_kernel_exec_info infoType, bool infoValue) { return defaultKernel->setUnifiedMemoryProperty(infoType, infoValue); }
+    void setSvmKernelExecInfo(GraphicsAllocation *argValue) { return defaultKernel->setSvmKernelExecInfo(argValue); }
+    void clearSvmKernelExecInfo() { return defaultKernel->clearSvmKernelExecInfo(); }
+    void setUnifiedMemoryExecInfo(GraphicsAllocation *argValue) { return defaultKernel->setUnifiedMemoryExecInfo(argValue); }
+    void clearUnifiedMemoryExecInfo() { return defaultKernel->clearUnifiedMemoryExecInfo(); }
+    int setKernelThreadArbitrationPolicy(uint32_t propertyValue) { return defaultKernel->setKernelThreadArbitrationPolicy(propertyValue); }
+    cl_int setKernelExecutionType(cl_execution_info_kernel_type_intel executionType) { return defaultKernel->setKernelExecutionType(executionType); }
+    int32_t setAdditionalKernelExecInfoWithParam(uint32_t paramName, size_t paramValueSize, const void *paramValue) { return defaultKernel->setAdditionalKernelExecInfoWithParam(paramName, paramValueSize, paramValue); }
 
   protected:
-    Kernel *kernel = nullptr;
+    KernelVectorType kernels;
+    Kernel *defaultKernel = nullptr;
 };
 
 } // namespace NEO
diff --git a/opencl/test/unit_test/accelerators/media_image_arg_tests.cpp b/opencl/test/unit_test/accelerators/media_image_arg_tests.cpp
index 222b7b8e91..a947ccd6a5 100644
--- a/opencl/test/unit_test/accelerators/media_image_arg_tests.cpp
+++ b/opencl/test/unit_test/accelerators/media_image_arg_tests.cpp
@@ -49,10 +49,11 @@ class MediaImageSetArgTest : public ClDeviceFixture,
         pKernelInfo->kernelArgInfo[1].isImage = true;
         pKernelInfo->kernelArgInfo[0].isImage = true;
 
-        pKernel = new MockKernel(program.get(), MockKernel::toKernelInfoContainer(*pKernelInfo, rootDeviceIndex));
+        int32_t retVal = CL_INVALID_PLATFORM;
+        pMultiDeviceKernel = MultiDeviceKernel::create<MockKernel>(program.get(), MockKernel::toKernelInfoContainer(*pKernelInfo, rootDeviceIndex), &retVal);
+        pKernel = static_cast<MockKernel *>(pMultiDeviceKernel->getKernel(rootDeviceIndex));
         ASSERT_NE(nullptr, pKernel);
-        ASSERT_EQ(CL_SUCCESS, pKernel->initialize());
-        pMultiDeviceKernel = new MultiDeviceKernel(pKernel);
+        ASSERT_EQ(CL_SUCCESS, retVal);
 
         ASSERT_EQ(true, pKernel->isVmeKernel());
 
diff --git a/opencl/test/unit_test/api/cl_api_tests.h b/opencl/test/unit_test/api/cl_api_tests.h
index 121de7b442..b62a381b5d 100644
--- a/opencl/test/unit_test/api/cl_api_tests.h
+++ b/opencl/test/unit_test/api/cl_api_tests.h
@@ -49,9 +49,9 @@ struct ApiFixture {
 
         pProgram = new MockProgram(pContext, false, toClDeviceVector(*pDevice));
 
-        pKernel = new MockKernel(pProgram, MockKernel::toKernelInfoContainer(pProgram->mockKernelInfo, testedRootDeviceIndex));
+        pMultiDeviceKernel = MockMultiDeviceKernel::create<MockKernel>(pProgram, MockKernel::toKernelInfoContainer(pProgram->mockKernelInfo, testedRootDeviceIndex));
+        pKernel = static_cast<MockKernel *>(pMultiDeviceKernel->getKernel(testedRootDeviceIndex));
         ASSERT_NE(nullptr, pKernel);
-        pMultiDeviceKernel = new MultiDeviceKernel(pKernel);
     }
 
     virtual void TearDown() {
diff --git a/opencl/test/unit_test/api/cl_get_kernel_max_concurrent_work_group_count_intel_tests.inl b/opencl/test/unit_test/api/cl_get_kernel_max_concurrent_work_group_count_intel_tests.inl
index b77a353948..5dfdda165a 100644
--- a/opencl/test/unit_test/api/cl_get_kernel_max_concurrent_work_group_count_intel_tests.inl
+++ b/opencl/test/unit_test/api/cl_get_kernel_max_concurrent_work_group_count_intel_tests.inl
@@ -72,7 +72,7 @@ TEST_F(clGetKernelMaxConcurrentWorkGroupCountTests, GivenVariousInputWhenGetting
     EXPECT_EQ(expectedMaxConcurrentWorkGroupCount, maxConcurrentWorkGroupCount);
 
     auto pKernelWithExecutionEnvironmentPatch = MockKernel::create(pCommandQueue->getDevice(), pProgram);
-    MultiDeviceKernel multiDeviceKernelWithExecutionEnvironmentPatch(pKernelWithExecutionEnvironmentPatch);
+    MultiDeviceKernel multiDeviceKernelWithExecutionEnvironmentPatch(MockMultiDeviceKernel::toKernelVector(pKernelWithExecutionEnvironmentPatch));
     retVal = clGetKernelMaxConcurrentWorkGroupCountINTEL(pCommandQueue, &multiDeviceKernelWithExecutionEnvironmentPatch, workDim,
                                                          globalWorkOffset, localWorkSize,
                                                          &maxConcurrentWorkGroupCount);
diff --git a/opencl/test/unit_test/api/cl_get_kernel_sub_group_info_khr_tests.inl b/opencl/test/unit_test/api/cl_get_kernel_sub_group_info_khr_tests.inl
index cefb9fa2f9..04745e800e 100644
--- a/opencl/test/unit_test/api/cl_get_kernel_sub_group_info_khr_tests.inl
+++ b/opencl/test/unit_test/api/cl_get_kernel_sub_group_info_khr_tests.inl
@@ -199,8 +199,8 @@ TEST_F(KernelSubGroupInfoKhrTest, GivenNullDeviceWhenGettingSubGroupInfoFromMult
 
     MockUnrestrictiveContext context;
     auto mockProgram = std::make_unique<MockProgram>(&context, false, context.getDevices());
-    auto mockKernel = new MockKernel(mockProgram.get(), pKernel->getKernelInfos());
-    auto pMultiDeviceKernel = std::make_unique<MultiDeviceKernel>(mockKernel);
+    std::unique_ptr<MultiDeviceKernel> pMultiDeviceKernel(
+        MultiDeviceKernel::create<MockKernel>(mockProgram.get(), pKernel->getKernelInfos(), nullptr));
 
     retVal = clGetKernelSubGroupInfoKHR(
         pMultiDeviceKernel.get(),
diff --git a/opencl/test/unit_test/api/cl_get_kernel_sub_group_info_tests.inl b/opencl/test/unit_test/api/cl_get_kernel_sub_group_info_tests.inl
index 97e5d0356d..f0ea6b20ca 100644
--- a/opencl/test/unit_test/api/cl_get_kernel_sub_group_info_tests.inl
+++ b/opencl/test/unit_test/api/cl_get_kernel_sub_group_info_tests.inl
@@ -348,8 +348,7 @@ TEST_F(KernelSubGroupInfoTest, GivenNullDeviceWhenGettingSubGroupInfoFromMultiDe
 
     MockUnrestrictiveContext context;
     auto mockProgram = std::make_unique<MockProgram>(&context, false, context.getDevices());
-    auto mockKernel = new MockKernel(mockProgram.get(), pKernel->getKernelInfos());
-    auto pMultiDeviceKernel = std::make_unique<MultiDeviceKernel>(mockKernel);
+    std::unique_ptr<MultiDeviceKernel> pMultiDeviceKernel(MultiDeviceKernel::create<MockKernel>(mockProgram.get(), pKernel->getKernelInfos(), nullptr));
 
     retVal = clGetKernelSubGroupInfo(
         pMultiDeviceKernel.get(),
diff --git a/opencl/test/unit_test/api/cl_get_kernel_suggested_local_work_size_intel_tests.inl b/opencl/test/unit_test/api/cl_get_kernel_suggested_local_work_size_intel_tests.inl
index bdfe725dc3..76328fc96a 100644
--- a/opencl/test/unit_test/api/cl_get_kernel_suggested_local_work_size_intel_tests.inl
+++ b/opencl/test/unit_test/api/cl_get_kernel_suggested_local_work_size_intel_tests.inl
@@ -98,7 +98,7 @@ TEST_F(clGetKernelSuggestedLocalWorkSizeTests, GivenVariousInputWhenGettingSugge
 
 TEST_F(clGetKernelSuggestedLocalWorkSizeTests, GivenKernelWithExecutionEnvironmentPatchedWhenGettingSuggestedLocalWorkSizeThenCorrectValuesAreReturned) {
     auto pKernelWithExecutionEnvironmentPatch = MockKernel::create(pCommandQueue->getDevice(), pProgram);
-    MultiDeviceKernel multiDeviceKernelWithExecutionEnvironmentPatch(pKernelWithExecutionEnvironmentPatch);
+    MultiDeviceKernel multiDeviceKernelWithExecutionEnvironmentPatch(MockMultiDeviceKernel::toKernelVector(pKernelWithExecutionEnvironmentPatch));
 
     size_t globalWorkOffset[] = {0, 0, 0};
     size_t globalWorkSize[] = {128, 128, 128};
diff --git a/opencl/test/unit_test/api/cl_get_kernel_work_group_info_tests.inl b/opencl/test/unit_test/api/cl_get_kernel_work_group_info_tests.inl
index f42adf20a5..7d485f5080 100644
--- a/opencl/test/unit_test/api/cl_get_kernel_work_group_info_tests.inl
+++ b/opencl/test/unit_test/api/cl_get_kernel_work_group_info_tests.inl
@@ -61,8 +61,8 @@ TEST_F(clGetKernelWorkGroupInfoTest, GivenNullDeviceWhenGettingWorkGroupInfoFrom
     size_t paramValueSizeRet;
     MockUnrestrictiveContext context;
     auto mockProgram = std::make_unique<MockProgram>(&context, false, context.getDevices());
-    auto pMockKernel = new MockKernel(mockProgram.get(), MockKernel::toKernelInfoContainer(pKernel->getKernelInfo(testedRootDeviceIndex), context.getDevice(0)->getRootDeviceIndex()));
-    auto pMultiDeviceKernel = std::make_unique<MultiDeviceKernel>(pMockKernel);
+    std::unique_ptr<MultiDeviceKernel> pMultiDeviceKernel(
+        MockMultiDeviceKernel::create<MockKernel>(mockProgram.get(), MockKernel::toKernelInfoContainer(pKernel->getKernelInfo(testedRootDeviceIndex), context.getDevice(0)->getRootDeviceIndex())));
 
     retVal = clGetKernelWorkGroupInfo(
         pMultiDeviceKernel.get(),
diff --git a/opencl/test/unit_test/api/cl_set_kernel_arg_svm_pointer_tests.inl b/opencl/test/unit_test/api/cl_set_kernel_arg_svm_pointer_tests.inl
index 6c3f586591..03b88890f0 100644
--- a/opencl/test/unit_test/api/cl_set_kernel_arg_svm_pointer_tests.inl
+++ b/opencl/test/unit_test/api/cl_set_kernel_arg_svm_pointer_tests.inl
@@ -40,9 +40,9 @@ class KernelArgSvmFixture : public ApiFixture<> {
         pKernelInfo->kernelArgInfo[0].kernelArgPatchInfoVector[0].size = (uint32_t)sizeof(void *);
         pKernelInfo->kernelArgInfo[0].metadata.addressQualifier = KernelArgMetadata::AddrGlobal;
 
-        pMockKernel = new MockKernel(pProgram, MockKernel::toKernelInfoContainer(*pKernelInfo, testedRootDeviceIndex));
-        ASSERT_EQ(CL_SUCCESS, pMockKernel->initialize());
-        pMockMultiDeviceKernel = new MultiDeviceKernel(pMockKernel);
+        pMockMultiDeviceKernel = MultiDeviceKernel::create<MockKernel>(pProgram, MockKernel::toKernelInfoContainer(*pKernelInfo, testedRootDeviceIndex), nullptr);
+        pMockKernel = static_cast<MockKernel *>(pMockMultiDeviceKernel->getKernel(testedRootDeviceIndex));
+        ASSERT_NE(nullptr, pMockKernel);
         pMockKernel->setCrossThreadData(pCrossThreadData, sizeof(pCrossThreadData));
     }
 
@@ -88,8 +88,9 @@ TEST_F(clSetKernelArgSVMPointerTests, GivenDeviceNotSupportingSvmWhenSettingKern
     auto hwInfo = executionEnvironment->rootDeviceEnvironments[ApiFixture::testedRootDeviceIndex]->getMutableHardwareInfo();
     hwInfo->capabilityTable.ftrSvm = false;
 
-    auto pMockKernel = new MockKernel(pProgram, MockKernel::toKernelInfoContainer(*pKernelInfo, testedRootDeviceIndex));
-    auto pMultiDeviceKernel = std::make_unique<MultiDeviceKernel>(pMockKernel);
+    std::unique_ptr<MultiDeviceKernel> pMultiDeviceKernel(
+        MultiDeviceKernel::create<MockKernel>(pProgram, MockKernel::toKernelInfoContainer(*pKernelInfo, testedRootDeviceIndex), nullptr));
+
     auto retVal = clSetKernelArgSVMPointer(
         pMultiDeviceKernel.get(), // cl_kernel kernel
         (cl_uint)-1,              // cl_uint arg_index
diff --git a/opencl/test/unit_test/api/cl_set_kernel_exec_info_tests.inl b/opencl/test/unit_test/api/cl_set_kernel_exec_info_tests.inl
index 0303d11dc1..d5cf00ec3f 100644
--- a/opencl/test/unit_test/api/cl_set_kernel_exec_info_tests.inl
+++ b/opencl/test/unit_test/api/cl_set_kernel_exec_info_tests.inl
@@ -22,9 +22,9 @@ class KernelExecInfoFixture : public ApiFixture<> {
         pKernelInfo = std::make_unique<KernelInfo>();
         pKernelInfo->kernelDescriptor.kernelAttributes.simdSize = 1;
 
-        pMockKernel = new MockKernel(pProgram, MockKernel::toKernelInfoContainer(*pKernelInfo, testedRootDeviceIndex));
-        ASSERT_EQ(CL_SUCCESS, pMockKernel->initialize());
-        pMockMultiDeviceKernel = new MultiDeviceKernel(pMockKernel);
+        pMockMultiDeviceKernel = MultiDeviceKernel::create<MockKernel>(pProgram, MockKernel::toKernelInfoContainer(*pKernelInfo, testedRootDeviceIndex), nullptr);
+        pMockKernel = static_cast<MockKernel *>(pMockMultiDeviceKernel->getKernel(testedRootDeviceIndex));
+        ASSERT_NE(nullptr, pMockKernel);
         svmCapabilities = pDevice->getDeviceInfo().svmCapabilities;
         if (svmCapabilities != 0) {
             ptrSvm = clSVMAlloc(pContext, CL_MEM_READ_WRITE, 256, 4);
@@ -70,8 +70,7 @@ TEST_F(clSetKernelArgSVMPointerTests, GivenDeviceNotSupportingSvmWhenSettingKern
     auto hwInfo = executionEnvironment->rootDeviceEnvironments[ApiFixture::testedRootDeviceIndex]->getMutableHardwareInfo();
     hwInfo->capabilityTable.ftrSvm = false;
 
-    auto pMockKernel = new MockKernel(pProgram, MockKernel::toKernelInfoContainer(*pKernelInfo, testedRootDeviceIndex));
-    auto pMultiDeviceKernel = std::make_unique<MultiDeviceKernel>(pMockKernel);
+    std::unique_ptr<MultiDeviceKernel> pMultiDeviceKernel(MultiDeviceKernel::create<MockKernel>(pProgram, MockKernel::toKernelInfoContainer(*pKernelInfo, testedRootDeviceIndex), nullptr));
     auto retVal = clSetKernelExecInfo(
         pMultiDeviceKernel.get(),     // cl_kernel kernel
         CL_KERNEL_EXEC_INFO_SVM_PTRS, // cl_kernel_exec_info param_name
diff --git a/opencl/test/unit_test/built_ins/built_in_tests.cpp b/opencl/test/unit_test/built_ins/built_in_tests.cpp
index af626473be..a7ca790981 100644
--- a/opencl/test/unit_test/built_ins/built_in_tests.cpp
+++ b/opencl/test/unit_test/built_ins/built_in_tests.cpp
@@ -497,7 +497,7 @@ HWTEST2_P(AuxBuiltInTests, givenKernelWithAuxTranslationRequiredWhenEnqueueCalle
 
     auto mockProgram = clUniquePtr(new MockProgram(toClDeviceVector(*pClDevice)));
     auto mockBuiltinKernel = MockKernel::create(*pDevice, mockProgram.get());
-    auto pMultiDeviceKernel = new MockMultiDeviceKernel(mockBuiltinKernel);
+    auto pMultiDeviceKernel = new MockMultiDeviceKernel(MockMultiDeviceKernel::toKernelVector(mockBuiltinKernel));
     mockAuxBuiltInOp->usedKernels.at(0).reset(pMultiDeviceKernel);
 
     MockKernelWithInternals mockKernel(*pClDevice, pContext);
diff --git a/opencl/test/unit_test/command_queue/enqueue_debug_kernel_tests.cpp b/opencl/test/unit_test/command_queue/enqueue_debug_kernel_tests.cpp
index bfa801ea36..5decccc349 100644
--- a/opencl/test/unit_test/command_queue/enqueue_debug_kernel_tests.cpp
+++ b/opencl/test/unit_test/command_queue/enqueue_debug_kernel_tests.cpp
@@ -51,14 +51,14 @@ class EnqueueDebugKernelTest : public ProgramSimpleFixture,
             ASSERT_EQ(CL_SUCCESS, retVal);
 
             // create a kernel
-            debugKernel = Kernel::create(
+            pMultiDeviceKernel = MultiDeviceKernel::create(
                 pProgram,
                 pProgram->getKernelInfosForKernel("CopyBuffer"),
                 &retVal);
+            debugKernel = pMultiDeviceKernel->getKernel(rootDeviceIndex);
 
             ASSERT_EQ(CL_SUCCESS, retVal);
             ASSERT_NE(nullptr, debugKernel);
-            pMultiDeviceKernel = new MultiDeviceKernel(debugKernel);
 
             cl_mem src = &bufferSrc;
             cl_mem dst = &bufferDst;
diff --git a/opencl/test/unit_test/command_queue/enqueue_kernel_1_tests.cpp b/opencl/test/unit_test/command_queue/enqueue_kernel_1_tests.cpp
index 7fbfd7a3bc..c3e3baf6d7 100644
--- a/opencl/test/unit_test/command_queue/enqueue_kernel_1_tests.cpp
+++ b/opencl/test/unit_test/command_queue/enqueue_kernel_1_tests.cpp
@@ -44,11 +44,11 @@ TEST_F(EnqueueKernelTest, givenKernelWhenAllArgsAreSetThenClEnqueueNDRangeKernel
     const size_t n = 512;
     size_t globalWorkSize[3] = {n, 1, 1};
     size_t localWorkSize[3] = {64, 1, 1};
-    cl_int retVal = CL_SUCCESS;
+    cl_int retVal = CL_INVALID_KERNEL;
     CommandQueue *pCmdQ2 = createCommandQueue(pClDevice);
 
-    auto kernel = Kernel::create(pProgram, pProgram->getKernelInfosForKernel("CopyBuffer"), &retVal);
-    auto pMultiDeviceKernel = std::make_unique<MultiDeviceKernel>(kernel);
+    std::unique_ptr<MultiDeviceKernel> pMultiDeviceKernel(MultiDeviceKernel::create(pProgram, pProgram->getKernelInfosForKernel("CopyBuffer"), &retVal));
+    auto kernel = pMultiDeviceKernel->getKernel(rootDeviceIndex);
 
     EXPECT_EQ(CL_SUCCESS, retVal);
 
@@ -88,8 +88,9 @@ TEST_F(EnqueueKernelTest, givenKernelWhenNotAllArgsAreSetButSetKernelArgIsCalled
     cl_int retVal = CL_SUCCESS;
     CommandQueue *pCmdQ2 = createCommandQueue(pClDevice);
 
-    auto kernel = Kernel::create(pProgram, pProgram->getKernelInfosForKernel("CopyBuffer"), &retVal);
-    auto pMultiDeviceKernel = std::make_unique<MultiDeviceKernel>(kernel);
+    std::unique_ptr<MultiDeviceKernel> pMultiDeviceKernel(MultiDeviceKernel::create(pProgram, pProgram->getKernelInfosForKernel("CopyBuffer"), &retVal));
+    auto kernel = pMultiDeviceKernel->getKernel(rootDeviceIndex);
+
     EXPECT_EQ(CL_SUCCESS, retVal);
 
     auto b0 = clCreateBuffer(context, 0, n * sizeof(float), nullptr, nullptr);
@@ -128,8 +129,9 @@ TEST_F(EnqueueKernelTest, givenKernelWhenSetKernelArgIsCalledForEachArgButAtLeas
     cl_int retVal = CL_SUCCESS;
     CommandQueue *pCmdQ2 = createCommandQueue(pClDevice);
 
-    auto kernel = Kernel::create(pProgram, pProgram->getKernelInfosForKernel("CopyBuffer"), &retVal);
-    auto pMultiDeviceKernel = std::make_unique<MultiDeviceKernel>(kernel);
+    std::unique_ptr<MultiDeviceKernel> pMultiDeviceKernel(MultiDeviceKernel::create(pProgram, pProgram->getKernelInfosForKernel("CopyBuffer"), &retVal));
+    auto kernel = pMultiDeviceKernel->getKernel(rootDeviceIndex);
+
     EXPECT_EQ(CL_SUCCESS, retVal);
 
     auto b0 = clCreateBuffer(context, 0, n * sizeof(float), nullptr, nullptr);
@@ -252,8 +254,9 @@ TEST_F(EnqueueKernelTest, givenKernelWhenAllArgsAreSetThenClEnqueueNDCountKernel
         pCmdQ2->getGpgpuEngine().osContext = pCmdQ2->getDevice().getEngine(aub_stream::ENGINE_CCS, EngineUsage::LowPriority).osContext;
     }
 
-    auto kernel = Kernel::create(pProgram, pProgram->getKernelInfosForKernel("CopyBuffer"), &retVal);
-    auto pMultiDeviceKernel = std::make_unique<MultiDeviceKernel>(kernel);
+    std::unique_ptr<MultiDeviceKernel> pMultiDeviceKernel(MultiDeviceKernel::create(pProgram, pProgram->getKernelInfosForKernel("CopyBuffer"), &retVal));
+    auto kernel = pMultiDeviceKernel->getKernel(rootDeviceIndex);
+
     EXPECT_EQ(CL_SUCCESS, retVal);
 
     auto b0 = clCreateBuffer(context, 0, n * sizeof(float), nullptr, nullptr);
@@ -297,8 +300,9 @@ TEST_F(EnqueueKernelTest, givenKernelWhenNotAllArgsAreSetButSetKernelArgIsCalled
         pCmdQ2->getGpgpuEngine().osContext = pCmdQ2->getDevice().getEngine(aub_stream::ENGINE_CCS, EngineUsage::LowPriority).osContext;
     }
 
-    auto kernel = Kernel::create(pProgram, pProgram->getKernelInfosForKernel("CopyBuffer"), &retVal);
-    auto pMultiDeviceKernel = std::make_unique<MultiDeviceKernel>(kernel);
+    std::unique_ptr<MultiDeviceKernel> pMultiDeviceKernel(MultiDeviceKernel::create(pProgram, pProgram->getKernelInfosForKernel("CopyBuffer"), &retVal));
+    auto kernel = pMultiDeviceKernel->getKernel(rootDeviceIndex);
+
     EXPECT_EQ(CL_SUCCESS, retVal);
 
     auto b0 = clCreateBuffer(context, 0, n * sizeof(float), nullptr, nullptr);
@@ -342,8 +346,9 @@ TEST_F(EnqueueKernelTest, givenKernelWhenSetKernelArgIsCalledForEachArgButAtLeas
         pCmdQ2->getGpgpuEngine().osContext = pCmdQ2->getDevice().getEngine(aub_stream::ENGINE_CCS, EngineUsage::LowPriority).osContext;
     }
 
-    auto kernel = Kernel::create(pProgram, pProgram->getKernelInfosForKernel("CopyBuffer"), &retVal);
-    auto pMultiDeviceKernel = std::make_unique<MultiDeviceKernel>(kernel);
+    std::unique_ptr<MultiDeviceKernel> pMultiDeviceKernel(MultiDeviceKernel::create(pProgram, pProgram->getKernelInfosForKernel("CopyBuffer"), &retVal));
+    auto kernel = pMultiDeviceKernel->getKernel(rootDeviceIndex);
+
     EXPECT_EQ(CL_SUCCESS, retVal);
 
     auto b0 = clCreateBuffer(context, 0, n * sizeof(float), nullptr, nullptr);
@@ -1268,8 +1273,9 @@ TEST_F(EnqueueKernelTest, givenKernelWhenAllArgsAreNotAndEventExistSetThenClEnqu
     cl_int retVal = CL_SUCCESS;
     CommandQueue *pCmdQ2 = createCommandQueue(pClDevice);
 
-    auto kernel = Kernel::create(pProgram, pProgram->getKernelInfosForKernel("CopyBuffer"), &retVal);
-    auto pMultiDeviceKernel = std::make_unique<MultiDeviceKernel>(kernel);
+    std::unique_ptr<MultiDeviceKernel> pMultiDeviceKernel(MultiDeviceKernel::create(pProgram, pProgram->getKernelInfosForKernel("CopyBuffer"), &retVal));
+    auto kernel = pMultiDeviceKernel->getKernel(rootDeviceIndex);
+
     EXPECT_EQ(CL_SUCCESS, retVal);
 
     EXPECT_FALSE(kernel->isPatched());
diff --git a/opencl/test/unit_test/command_queue/enqueue_svm_tests.cpp b/opencl/test/unit_test/command_queue/enqueue_svm_tests.cpp
index 3583f03ed3..fd98a50728 100644
--- a/opencl/test/unit_test/command_queue/enqueue_svm_tests.cpp
+++ b/opencl/test/unit_test/command_queue/enqueue_svm_tests.cpp
@@ -766,8 +766,8 @@ TEST_F(EnqueueSvmTest, givenEnqueueTaskBlockedOnUserEventWhenItIsEnqueuedThenSur
 
     auto program = clUniquePtr(Program::createBuiltInFromSource<MockProgram>("FillBufferBytes", context, context->getDevices(), &retVal));
     program->build(program->getDevices(), nullptr, false);
-    auto kernel = Kernel::create<MockKernel>(program.get(), program->getKernelInfosForKernel("FillBufferBytes"), &retVal);
-    MultiDeviceKernel multiDeviceKernel(kernel);
+    auto pMultiDeviceKernel = clUniquePtr(MultiDeviceKernel::create<MockKernel>(program.get(), program->getKernelInfosForKernel("FillBufferBytes"), &retVal));
+    auto kernel = static_cast<MockKernel *>(pMultiDeviceKernel->getKernel(rootDeviceIndex));
     std::vector<Surface *> allSurfaces;
     kernel->getResidency(allSurfaces, rootDeviceIndex);
     EXPECT_EQ(1u, allSurfaces.size());
diff --git a/opencl/test/unit_test/command_stream/command_stream_receiver_flush_task_3_tests.cpp b/opencl/test/unit_test/command_stream/command_stream_receiver_flush_task_3_tests.cpp
index 1b2252bf29..75882c69c2 100644
--- a/opencl/test/unit_test/command_stream/command_stream_receiver_flush_task_3_tests.cpp
+++ b/opencl/test/unit_test/command_stream/command_stream_receiver_flush_task_3_tests.cpp
@@ -1628,7 +1628,7 @@ HWTEST_F(CommandStreamReceiverFlushTaskTests, GivenBlockedKernelWhenItIsUnblocke
     auto mockProgram = std::make_unique<MockProgram>(&mockContext, false, toClDeviceVector(*pClDevice));
 
     auto pKernel = MockKernel::create(*pDevice, mockProgram.get(), numGrfRequired);
-    MultiDeviceKernel multiDeviceKernel(pKernel);
+    MultiDeviceKernel multiDeviceKernel(MockMultiDeviceKernel::toKernelVector(pKernel));
     auto event = std::make_unique<MockEvent<Event>>(pCmdQ.get(), CL_COMMAND_MARKER, 0, 0);
     auto cmdStream = new LinearStream(pDevice->getMemoryManager()->allocateGraphicsMemoryWithProperties({pDevice->getRootDeviceIndex(), 4096, GraphicsAllocation::AllocationType::COMMAND_BUFFER, pDevice->getDeviceBitfield()}));
 
diff --git a/opencl/test/unit_test/execution_model/submit_blocked_parent_kernel_tests.cpp b/opencl/test/unit_test/execution_model/submit_blocked_parent_kernel_tests.cpp
index c7e116f1ec..d34831db42 100644
--- a/opencl/test/unit_test/execution_model/submit_blocked_parent_kernel_tests.cpp
+++ b/opencl/test/unit_test/execution_model/submit_blocked_parent_kernel_tests.cpp
@@ -81,7 +81,7 @@ HWCMDTEST_F(IGFX_GEN8_CORE, ParentKernelCommandQueueFixture, givenLockedEMcritca
 
     cl_queue_properties properties[3] = {0};
     MockParentKernel *parentKernel = MockParentKernel::create(*context);
-    MultiDeviceKernel multiDeviceKernel(parentKernel);
+    MultiDeviceKernel multiDeviceKernel(MockMultiDeviceKernel::toKernelVector(parentKernel));
     MockDeviceQueueHwWithCriticalSectionRelease<FamilyType> mockDevQueue(context, device, properties[0]);
     parentKernel->createReflectionSurface();
     context->setDefaultDeviceQueue(&mockDevQueue);
@@ -121,7 +121,7 @@ HWCMDTEST_F(IGFX_GEN8_CORE, ParentKernelCommandQueueFixture, givenParentKernelWh
 
     cl_queue_properties properties[3] = {0};
     MockParentKernel *parentKernel = MockParentKernel::create(*context);
-    MultiDeviceKernel multiDeviceKernel(parentKernel);
+    MultiDeviceKernel multiDeviceKernel(MockMultiDeviceKernel::toKernelVector(parentKernel));
     MockDeviceQueueHwWithCriticalSectionRelease<FamilyType> mockDevQueue(context, device, properties[0]);
     parentKernel->createReflectionSurface();
     context->setDefaultDeviceQueue(&mockDevQueue);
@@ -184,7 +184,7 @@ HWCMDTEST_F(IGFX_GEN8_CORE, ParentKernelCommandQueueFixture, givenParentKernelWh
 
     cl_queue_properties properties[3] = {0};
     MockParentKernel *parentKernel = MockParentKernel::create(*context);
-    MultiDeviceKernel multiDeviceKernel(parentKernel);
+    MultiDeviceKernel multiDeviceKernel(MockMultiDeviceKernel::toKernelVector(parentKernel));
     MockDeviceQueueHwWithCriticalSectionRelease<FamilyType> mockDevQueue(context, device, properties[0]);
     parentKernel->createReflectionSurface();
     context->setDefaultDeviceQueue(&mockDevQueue);
@@ -224,7 +224,7 @@ HWCMDTEST_F(IGFX_GEN8_CORE, ParentKernelCommandQueueFixture, givenBlockedParentK
 
     cl_queue_properties properties[3] = {0};
     MockParentKernel *parentKernel = MockParentKernel::create(*context);
-    MultiDeviceKernel multiDeviceKernel(parentKernel);
+    MultiDeviceKernel multiDeviceKernel(MockMultiDeviceKernel::toKernelVector(parentKernel));
     MockDeviceQueueHwWithCriticalSectionRelease<FamilyType> mockDevQueue(context, device, properties[0]);
     parentKernel->createReflectionSurface();
     context->setDefaultDeviceQueue(&mockDevQueue);
@@ -264,7 +264,7 @@ HWCMDTEST_F(IGFX_GEN8_CORE, ParentKernelCommandQueueFixture, givenParentKernelWh
 
     cl_queue_properties properties[3] = {0};
     MockParentKernel *parentKernel = MockParentKernel::create(*context);
-    MultiDeviceKernel multiDeviceKernel(parentKernel);
+    MultiDeviceKernel multiDeviceKernel(MockMultiDeviceKernel::toKernelVector(parentKernel));
     MockDeviceQueueHwWithCriticalSectionRelease<FamilyType> mockDevQueue(context, device, properties[0]);
     parentKernel->createReflectionSurface();
     context->setDefaultDeviceQueue(&mockDevQueue);
@@ -302,7 +302,7 @@ HWCMDTEST_F(IGFX_GEN8_CORE, ParentKernelCommandQueueFixture, givenUsedCommandQue
 
     cl_queue_properties properties[3] = {0};
     MockParentKernel *parentKernel = MockParentKernel::create(*context);
-    MultiDeviceKernel multiDeviceKernel(parentKernel);
+    MultiDeviceKernel multiDeviceKernel(MockMultiDeviceKernel::toKernelVector(parentKernel));
     MockDeviceQueueHw<FamilyType> mockDevQueue(context, device, properties[0]);
     parentKernel->createReflectionSurface();
     context->setDefaultDeviceQueue(&mockDevQueue);
@@ -358,7 +358,7 @@ HWCMDTEST_F(IGFX_GEN8_CORE, ParentKernelCommandQueueFixture, givenNotUsedSSHWhen
 
     cl_queue_properties properties[3] = {0};
     MockParentKernel *parentKernel = MockParentKernel::create(*context);
-    MultiDeviceKernel multiDeviceKernel(parentKernel);
+    MultiDeviceKernel multiDeviceKernel(MockMultiDeviceKernel::toKernelVector(parentKernel));
     MockDeviceQueueHw<FamilyType> mockDevQueue(context, device, properties[0]);
     parentKernel->createReflectionSurface();
     context->setDefaultDeviceQueue(&mockDevQueue);
@@ -403,7 +403,7 @@ HWCMDTEST_F(IGFX_GEN8_CORE, ParentKernelCommandQueueFixture, givenBlockedCommand
 
     cl_queue_properties properties[3] = {0};
     auto parentKernel = MockParentKernel::create(*context);
-    MultiDeviceKernel multiDeviceKernel(parentKernel);
+    MultiDeviceKernel multiDeviceKernel(MockMultiDeviceKernel::toKernelVector(parentKernel));
 
     MockDeviceQueueHw<FamilyType> mockDevQueue(context, device, properties[0]);
     parentKernel->createReflectionSurface();
diff --git a/opencl/test/unit_test/fixtures/execution_model_fixture.h b/opencl/test/unit_test/fixtures/execution_model_fixture.h
index 6c640a8d53..91175f94cc 100644
--- a/opencl/test/unit_test/fixtures/execution_model_fixture.h
+++ b/opencl/test/unit_test/fixtures/execution_model_fixture.h
@@ -93,7 +93,7 @@ class ExecutionModelSchedulerTest : public ClDeviceFixture,
 
         parentKernel = MockParentKernel::create(*context);
         ASSERT_NE(nullptr, parentKernel);
-        pMultiDeviceKernel = new MockMultiDeviceKernel(parentKernel);
+        pMultiDeviceKernel = new MockMultiDeviceKernel(MockMultiDeviceKernel::toKernelVector(parentKernel));
     }
 
     void TearDown() override {
diff --git a/opencl/test/unit_test/fixtures/execution_model_kernel_fixture.h b/opencl/test/unit_test/fixtures/execution_model_kernel_fixture.h
index c6bcf77e51..421928bab1 100644
--- a/opencl/test/unit_test/fixtures/execution_model_kernel_fixture.h
+++ b/opencl/test/unit_test/fixtures/execution_model_kernel_fixture.h
@@ -43,7 +43,7 @@ struct ExecutionModelKernelFixture : public ProgramFromBinaryFixture,
             pProgram->getKernelInfosForKernel(kernelName),
             &retVal);
 
-        pMultiDeviceKernel = new MockMultiDeviceKernel(pKernel);
+        pMultiDeviceKernel = new MockMultiDeviceKernel(MockMultiDeviceKernel::toKernelVector(pKernel));
         ASSERT_EQ(CL_SUCCESS, retVal);
         ASSERT_NE(nullptr, pKernel);
     }
diff --git a/opencl/test/unit_test/fixtures/kernel_arg_fixture.cpp b/opencl/test/unit_test/fixtures/kernel_arg_fixture.cpp
index a1fd508e8e..7702434f1c 100644
--- a/opencl/test/unit_test/fixtures/kernel_arg_fixture.cpp
+++ b/opencl/test/unit_test/fixtures/kernel_arg_fixture.cpp
@@ -63,9 +63,10 @@ void KernelImageArgTest::SetUp() {
     ClDeviceFixture::SetUp();
     context.reset(new MockContext(pClDevice));
     program = std::make_unique<MockProgram>(context.get(), false, toClDeviceVector(*pClDevice));
-    pKernel = new MockKernel(program.get(), MockKernel::toKernelInfoContainer(*pKernelInfo, rootDeviceIndex));
-    ASSERT_EQ(CL_SUCCESS, pKernel->initialize());
-    pMultiDeviceKernel = std::make_unique<MultiDeviceKernel>(pKernel);
+    int32_t retVal = CL_INVALID_VALUE;
+    pMultiDeviceKernel.reset(MultiDeviceKernel::create<MockKernel>(program.get(), MockKernel::toKernelInfoContainer(*pKernelInfo, rootDeviceIndex), &retVal));
+    pKernel = static_cast<MockKernel *>(pMultiDeviceKernel->getKernel(rootDeviceIndex));
+    ASSERT_EQ(CL_SUCCESS, retVal);
 
     pKernel->setKernelArgHandler(0, &Kernel::setArgImage);
     pKernel->setKernelArgHandler(1, &Kernel::setArgImage);
diff --git a/opencl/test/unit_test/fixtures/media_kernel_fixture.h b/opencl/test/unit_test/fixtures/media_kernel_fixture.h
index 2809817ac6..e55ef3b633 100644
--- a/opencl/test/unit_test/fixtures/media_kernel_fixture.h
+++ b/opencl/test/unit_test/fixtures/media_kernel_fixture.h
@@ -81,7 +81,7 @@ struct MediaKernelFixture : public HelloWorldFixture<FactoryType>,
 
         ASSERT_NE(nullptr, pVmeKernel);
         ASSERT_EQ(true, pVmeKernel->isVmeKernel());
-        pMultiDeviceVmeKernel = new MockMultiDeviceKernel(pVmeKernel);
+        pMultiDeviceVmeKernel = new MockMultiDeviceKernel(MockMultiDeviceKernel::toKernelVector(pVmeKernel));
     }
 
     void TearDown() override {
diff --git a/opencl/test/unit_test/gen12lp/gpgpu_walker_tests_gen12lp.cpp b/opencl/test/unit_test/gen12lp/gpgpu_walker_tests_gen12lp.cpp
index a0d8cf59ae..4f089aa3f6 100644
--- a/opencl/test/unit_test/gen12lp/gpgpu_walker_tests_gen12lp.cpp
+++ b/opencl/test/unit_test/gen12lp/gpgpu_walker_tests_gen12lp.cpp
@@ -37,7 +37,7 @@ GEN12LPTEST_F(GpgpuWalkerTests, givenMiStoreRegMemWhenAdjustMiStoreRegMemModeThe
 
 class MockKernelWithApplicableWa : public MockKernel {
   public:
-    MockKernelWithApplicableWa(Program *program, KernelInfoContainer &kernelInfos) : MockKernel(program, kernelInfos) {}
+    MockKernelWithApplicableWa(Program *program, const KernelInfoContainer &kernelInfos) : MockKernel(program, kernelInfos) {}
     bool requiresWaDisableRccRhwoOptimization(uint32_t rootDeviceIndex) const override {
         return waApplicable;
     }
@@ -53,8 +53,8 @@ struct HardwareInterfaceTests : public ClDeviceFixture, public LinearStreamFixtu
         pCommandQueue = new MockCommandQueue(pContext, pClDevice, nullptr);
         pProgram = new MockProgram(pContext, false, toClDeviceVector(*pClDevice));
         auto kernelInfos = MockKernel::toKernelInfoContainer(pProgram->mockKernelInfo, rootDeviceIndex);
-        pKernel = new MockKernelWithApplicableWa(pProgram, kernelInfos);
-        pMultiDeviceKernel = new MultiDeviceKernel(pKernel);
+        pMultiDeviceKernel = MockMultiDeviceKernel::create<MockKernelWithApplicableWa>(pProgram, kernelInfos);
+        pKernel = static_cast<MockKernelWithApplicableWa *>(pMultiDeviceKernel->getKernel(rootDeviceIndex));
     }
 
     void TearDown() override {
diff --git a/opencl/test/unit_test/gtpin/gtpin_tests.cpp b/opencl/test/unit_test/gtpin/gtpin_tests.cpp
index adc1483c48..42627bbb37 100644
--- a/opencl/test/unit_test/gtpin/gtpin_tests.cpp
+++ b/opencl/test/unit_test/gtpin/gtpin_tests.cpp
@@ -1187,7 +1187,7 @@ HWCMDTEST_F(IGFX_GEN8_CORE, GTPinTests, givenInitializedGTPinInterfaceWhenKernel
     size_t localWorkSize[3] = {1, 1, 1};
 
     MockParentKernel *parentKernel = MockParentKernel::create(*pContext);
-    auto pMultiDeviceKernel = std::make_unique<MultiDeviceKernel>(parentKernel);
+    auto pMultiDeviceKernel = std::make_unique<MultiDeviceKernel>(MockMultiDeviceKernel::toKernelVector(parentKernel));
 
     retVal = clEnqueueNDRangeKernel(cmdQ, pMultiDeviceKernel.get(), workDim, globalWorkOffset, globalWorkSize, localWorkSize, 0, nullptr, nullptr);
     EXPECT_EQ(CL_SUCCESS, retVal);
@@ -2406,8 +2406,8 @@ TEST_F(GTPinTests, givenInitializedGTPinInterfaceWhenOnKernelSubitIsCalledThenCo
 
     auto pProgramm = std::make_unique<MockProgram>(context.get(), false, toClDeviceVector(*pDevice));
     std::unique_ptr<MockCommandQueue> cmdQ(new MockCommandQueue(context.get(), pDevice, nullptr));
-    auto pKernel = new MockKernel(pProgramm.get(), MockKernel::toKernelInfoContainer(*pKernelInfo, rootDeviceIndex));
-    auto pMultiDeviceKernel = std::make_unique<MultiDeviceKernel>(pKernel);
+    std::unique_ptr<MultiDeviceKernel> pMultiDeviceKernel(MockMultiDeviceKernel::create<MockKernel>(pProgramm.get(), MockKernel::toKernelInfoContainer(*pKernelInfo, rootDeviceIndex)));
+    auto pKernel = static_cast<MockKernel *>(pMultiDeviceKernel->getKernel(rootDeviceIndex));
 
     pKernel->setSshLocal(nullptr, sizeof(surfaceStateHeap), rootDeviceIndex);
 
diff --git a/opencl/test/unit_test/mem_obj/buffer_set_arg_tests.cpp b/opencl/test/unit_test/mem_obj/buffer_set_arg_tests.cpp
index c82a96d531..1c8b149e7d 100644
--- a/opencl/test/unit_test/mem_obj/buffer_set_arg_tests.cpp
+++ b/opencl/test/unit_test/mem_obj/buffer_set_arg_tests.cpp
@@ -67,10 +67,11 @@ class BufferSetArgTest : public ContextFixture,
 
         pProgram = new MockProgram(pContext, false, toClDeviceVector(*pClDevice));
 
-        pKernel = new MockKernel(pProgram, MockKernel::toKernelInfoContainer(*pKernelInfo, rootDeviceIndex));
+        retVal = CL_INVALID_VALUE;
+        pMultiDeviceKernel = MultiDeviceKernel::create<MockKernel>(pProgram, MockKernel::toKernelInfoContainer(*pKernelInfo, rootDeviceIndex), &retVal);
+        pKernel = static_cast<MockKernel *>(pMultiDeviceKernel->getKernel(rootDeviceIndex));
         ASSERT_NE(nullptr, pKernel);
-        ASSERT_EQ(CL_SUCCESS, pKernel->initialize());
-        pMultiDeviceKernel = new MultiDeviceKernel(pKernel);
+        ASSERT_EQ(CL_SUCCESS, retVal);
         pKernel->setCrossThreadData(pCrossThreadData, sizeof(pCrossThreadData));
 
         pKernel->setKernelArgHandler(1, &Kernel::setArgBuffer);
diff --git a/opencl/test/unit_test/mem_obj/image_set_arg_tests.cpp b/opencl/test/unit_test/mem_obj/image_set_arg_tests.cpp
index 4ac0817aa3..c8835cc648 100644
--- a/opencl/test/unit_test/mem_obj/image_set_arg_tests.cpp
+++ b/opencl/test/unit_test/mem_obj/image_set_arg_tests.cpp
@@ -80,10 +80,11 @@ class ImageSetArgTest : public ClDeviceFixture,
         pKernelInfo->kernelArgInfo[0].isImage = true;
 
         program = std::make_unique<MockProgram>(toClDeviceVector(*pClDevice));
-        pKernel = new MockKernel(program.get(), MockKernel::toKernelInfoContainer(*pKernelInfo, rootDeviceIndex));
+        retVal = CL_INVALID_VALUE;
+        pMultiDeviceKernel = MultiDeviceKernel::create<MockKernel>(program.get(), MockKernel::toKernelInfoContainer(*pKernelInfo, rootDeviceIndex), &retVal);
+        pKernel = static_cast<MockKernel *>(pMultiDeviceKernel->getKernel(rootDeviceIndex));
         ASSERT_NE(nullptr, pKernel);
-        ASSERT_EQ(CL_SUCCESS, pKernel->initialize());
-        pMultiDeviceKernel = new MultiDeviceKernel(pKernel);
+        ASSERT_EQ(CL_SUCCESS, retVal);
 
         pKernel->setKernelArgHandler(0, &Kernel::setArgImage);
         pKernel->setKernelArgHandler(1, &Kernel::setArgImage);
@@ -942,10 +943,11 @@ class ImageMediaBlockSetArgTest : public ImageSetArgTest {
         pKernelInfo->kernelArgInfo[0].isMediaBlockImage = true;
 
         program = std::make_unique<MockProgram>(toClDeviceVector(*pClDevice));
-        pKernel = new MockKernel(program.get(), MockKernel::toKernelInfoContainer(*pKernelInfo, rootDeviceIndex));
+        retVal = CL_INVALID_VALUE;
+        pMultiDeviceKernel = MultiDeviceKernel::create<MockKernel>(program.get(), MockKernel::toKernelInfoContainer(*pKernelInfo, rootDeviceIndex), &retVal);
+        pKernel = static_cast<MockKernel *>(pMultiDeviceKernel->getKernel(rootDeviceIndex));
         ASSERT_NE(nullptr, pKernel);
-        ASSERT_EQ(CL_SUCCESS, pKernel->initialize());
-        pMultiDeviceKernel = new MultiDeviceKernel(pKernel);
+        ASSERT_EQ(CL_SUCCESS, retVal);
 
         pKernel->setKernelArgHandler(0, &Kernel::setArgImage);
         pKernel->setKernelArgHandler(1, &Kernel::setArgImage);
diff --git a/opencl/test/unit_test/mocks/mock_kernel.h b/opencl/test/unit_test/mocks/mock_kernel.h
index c4dccdd693..b08c8f339f 100644
--- a/opencl/test/unit_test/mocks/mock_kernel.h
+++ b/opencl/test/unit_test/mocks/mock_kernel.h
@@ -73,8 +73,26 @@ struct MockKernelObjForAuxTranslation : public KernelObjForAuxTranslation {
 
 class MockMultiDeviceKernel : public MultiDeviceKernel {
   public:
+    static KernelVectorType toKernelVector(Kernel *pKernel) {
+        KernelVectorType kernelVector;
+        kernelVector.resize(pKernel->getProgram()->getMaxRootDeviceIndex() + 1);
+        kernelVector[pKernel->getProgram()->getDevices()[0]->getRootDeviceIndex()] = pKernel;
+        return kernelVector;
+    }
     using MultiDeviceKernel::MultiDeviceKernel;
-
+    template <typename kernel_t = Kernel>
+    static MockMultiDeviceKernel *create(Program *programArg, const KernelInfoContainer &kernelInfoArg) {
+        KernelVectorType kernelVector;
+        kernelVector.resize(programArg->getMaxRootDeviceIndex() + 1);
+        for (auto &pDevice : programArg->getDevices()) {
+            auto rootDeviceIndex = pDevice->getRootDeviceIndex();
+            if (kernelVector[rootDeviceIndex]) {
+                continue;
+            }
+            kernelVector[rootDeviceIndex] = new kernel_t(programArg, kernelInfoArg);
+        }
+        return new MockMultiDeviceKernel(std::move(kernelVector));
+    }
     void takeOwnership() const override {
         MultiDeviceKernel::takeOwnership();
         takeOwnershipCalls++;
@@ -367,7 +385,15 @@ class MockKernelWithInternals {
         mockProgram = new MockProgram(context, false, deviceVector);
         mockKernel = new MockKernel(mockProgram, kernelInfos);
         mockKernel->setCrossThreadData(&crossThreadData, sizeof(crossThreadData));
-        mockMultiDeviceKernel = new MockMultiDeviceKernel(mockKernel);
+        KernelVectorType mockKernels;
+        mockKernels.resize(mockProgram->getMaxRootDeviceIndex() + 1);
+        for (const auto &pClDevice : deviceVector) {
+            auto rootDeviceIndex = pClDevice->getRootDeviceIndex();
+            if (mockKernels[rootDeviceIndex] == nullptr) {
+                mockKernels[rootDeviceIndex] = mockKernel;
+            }
+        }
+        mockMultiDeviceKernel = new MockMultiDeviceKernel(std::move(mockKernels));
 
         for (const auto &pClDevice : deviceVector) {
             mockKernel->setSshLocal(&sshLocal, sizeof(sshLocal), pClDevice->getRootDeviceIndex());
diff --git a/opencl/test/unit_test/sampler/sampler_set_arg_tests.cpp b/opencl/test/unit_test/sampler/sampler_set_arg_tests.cpp
index 632374c912..b83ec6a86f 100644
--- a/opencl/test/unit_test/sampler/sampler_set_arg_tests.cpp
+++ b/opencl/test/unit_test/sampler/sampler_set_arg_tests.cpp
@@ -56,10 +56,11 @@ class SamplerSetArgFixture : public ClDeviceFixture {
         pKernelInfo->kernelArgInfo[1].isSampler = true;
 
         program = std::make_unique<MockProgram>(toClDeviceVector(*pClDevice));
-        pKernel = new MockKernel(program.get(), MockKernel::toKernelInfoContainer(*pKernelInfo, rootDeviceIndex));
+        retVal = CL_INVALID_VALUE;
+        pMultiDeviceKernel = MultiDeviceKernel::create<MockKernel>(program.get(), MockKernel::toKernelInfoContainer(*pKernelInfo, rootDeviceIndex), &retVal);
+        pKernel = static_cast<MockKernel *>(pMultiDeviceKernel->getKernel(rootDeviceIndex));
         ASSERT_NE(nullptr, pKernel);
-        ASSERT_EQ(CL_SUCCESS, pKernel->initialize());
-        pMultiDeviceKernel = new MultiDeviceKernel(pKernel);
+        ASSERT_EQ(CL_SUCCESS, retVal);
 
         pKernel->setKernelArgHandler(0, &Kernel::setArgSampler);
         pKernel->setKernelArgHandler(1, &Kernel::setArgSampler);