diff --git a/sycl/include/CL/sycl/detail/pi.def b/sycl/include/CL/sycl/detail/pi.def
index 2574cc3c4d485..e54f3c009dbfb 100644
--- a/sycl/include/CL/sycl/detail/pi.def
+++ b/sycl/include/CL/sycl/detail/pi.def
@@ -18,6 +18,7 @@
 _PI_API(piPlatformsGet)
 _PI_API(piPlatformGetInfo)
 // Device
+_PI_API(piextDeviceConvert)
 _PI_API(piDevicesGet)
 _PI_API(piDeviceGetInfo)
 _PI_API(piDevicePartition)
@@ -45,6 +46,7 @@ _PI_API(piMemRetain)
 _PI_API(piMemRelease)
 _PI_API(piMemBufferPartition)
 // Program
+_PI_API(piextProgramConvert)
 _PI_API(piProgramCreate)
 _PI_API(piclProgramCreateWithSource)
 _PI_API(piclProgramCreateWithBinary)
diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h
index 29533d67a767a..507f15fbd03ad 100644
--- a/sycl/include/CL/sycl/detail/pi.h
+++ b/sycl/include/CL/sycl/detail/pi.h
@@ -711,6 +711,16 @@ pi_result piPlatformGetInfo(pi_platform platform, pi_platform_info param_name,
 //
 // Device
 //
+///
+/// Create PI device from the given raw device handle (if the "device"
+/// points to null), or, vice versa, extract the raw device handle into
+/// the "handle" (if it was pointing to a null) from the given PI device.
+/// NOTE: The instance of the PI device created is retained.
+///
+pi_result piextDeviceConvert(
+    pi_device *device, ///< [in,out] the pointer to PI device
+    void **handle);    ///< [in,out] the pointer to the raw device handle
+
 pi_result piDevicesGet(pi_platform platform, pi_device_type device_type,
                        pi_uint32 num_entries, pi_device *devices,
                        pi_uint32 *num_devices);
@@ -811,6 +821,17 @@ pi_result piMemBufferPartition(pi_mem buffer, pi_mem_flags flags,
 //
 // Program
 //
+///
+/// Create PI program from the given raw program handle (if the "program"
+/// points to null), or, vice versa, extract the raw program handle into
+/// the "handle" (if it was pointing to a null) from the given PI program.
+/// NOTE: The instance of the PI program created is retained.
+///
+pi_result piextProgramConvert(
+    pi_context context,  ///< [in] the PI context of the program
+    pi_program *program, ///< [in,out] the pointer to PI program
+    void **handle);      ///< [in,out] the pointer to the raw program handle
+
 pi_result piProgramCreate(pi_context context, const void *il, size_t length,
                           pi_program *res_program);
 
diff --git a/sycl/include/CL/sycl/detail/pi.hpp b/sycl/include/CL/sycl/detail/pi.hpp
index f5404a4afc9a1..954e9d0c6ce20 100644
--- a/sycl/include/CL/sycl/detail/pi.hpp
+++ b/sycl/include/CL/sycl/detail/pi.hpp
@@ -171,12 +171,23 @@ namespace RT = cl::sycl::detail::pi;
 
 // Want all the needed casts be explicit, do not define conversion
 // operators.
-template <class To, class From> To pi::cast(From value) {
+template <class To, class From> To inline pi::cast(From value) {
   // TODO: see if more sanity checks are possible.
   RT::assertion((sizeof(From) == sizeof(To)), "assert: cast failed size check");
   return (To)(value);
 }
 
+// These conversions should use PI interop API.
+template <> pi::PiProgram inline pi::cast(cl_program interop) {
+  RT::assertion(false, "pi::cast -> use piextProgramConvert");
+  return {};
+}
+
+template <> pi::PiDevice inline pi::cast(cl_device_id interop) {
+  RT::assertion(false, "pi::cast -> use piextDeviceConvert");
+  return {};
+}
+
 } // namespace detail
 
 // For shortness of using PI from the top-level sycl files.
diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp
index 52d87d6818756..546da5a634fa6 100644
--- a/sycl/plugins/cuda/pi_cuda.cpp
+++ b/sycl/plugins/cuda/pi_cuda.cpp
@@ -680,6 +680,11 @@ pi_result cuda_piPlatformGetInfo(pi_platform platform,
   return {};
 }
 
+pi_result cuda_piextDeviceConvert(pi_device *device, void **handle) {
+  cl::sycl::detail::pi::die("cuda_piextDeviceConvert not implemented");
+  return {};
+}
+
 pi_result cuda_piDevicesGet(pi_platform platform, pi_device_type device_type,
                             pi_uint32 num_entries, pi_device *devices,
                             pi_uint32 *num_devices) {
@@ -2138,6 +2143,15 @@ pi_result cuda_piMemRetain(pi_mem mem) {
 //
 // Program
 //
+pi_result cuda_piextProgramConvert(
+    pi_context context,  ///< [in] the PI context of the program
+    pi_program *program, ///< [in,out] the pointer to PI program
+    void **handle)       ///< [in,out] the pointer to the raw program handle
+{
+  cl::sycl::detail::pi::die("cuda_piextProgramConvert not implemented");
+  return {};
+}
+
 pi_result cuda_piProgramCreate(pi_context context, const void *il,
                                size_t length, pi_program *res_program) {
   cl::sycl::detail::pi::die("cuda_piProgramCreate not implemented");
@@ -3480,6 +3494,7 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piPlatformsGet, cuda_piPlatformsGet)
   _PI_CL(piPlatformGetInfo, cuda_piPlatformGetInfo)
   // Device
+  _PI_CL(piextDeviceConvert, cuda_piextDeviceConvert)
   _PI_CL(piDevicesGet, cuda_piDevicesGet)
   _PI_CL(piDeviceGetInfo, cuda_piDeviceGetInfo)
   _PI_CL(piDevicePartition, cuda_piDevicePartition)
@@ -3507,6 +3522,7 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piMemRelease, cuda_piMemRelease)
   _PI_CL(piMemBufferPartition, cuda_piMemBufferPartition)
   // Program
+  _PI_CL(piextProgramConvert, cuda_piextProgramConvert)
   _PI_CL(piProgramCreate, cuda_piProgramCreate)
   _PI_CL(piclProgramCreateWithSource, cuda_piclProgramCreateWithSource)
   _PI_CL(piclProgramCreateWithBinary, cuda_piclProgramCreateWithBinary)
diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp
index 846d2b60a3708..51561e780c61b 100644
--- a/sycl/plugins/opencl/pi_opencl.cpp
+++ b/sycl/plugins/opencl/pi_opencl.cpp
@@ -176,6 +176,24 @@ pi_result OCL(piPlatformsGet)(pi_uint32 num_entries, pi_platform *platforms,
   return static_cast<pi_result>(result);
 }
 
+pi_result OCL(piextDeviceConvert)(pi_device *device, void **handle) {
+  // The PI device is the same as OpenCL device handle.
+  assert(device);
+  assert(handle);
+
+  if (*device == nullptr) {
+    // unitialized *device.
+    assert(*handle);
+    *device = cast<pi_device>(*handle);
+  } else {
+    assert(*handle == nullptr);
+    *handle = *device;
+  }
+
+  cl_int result = clRetainDevice(cast<cl_device_id>(*handle));
+  return cast<pi_result>(result);
+}
+
 // Example of a PI interface that does not map exactly to an OpenCL one.
 pi_result OCL(piDevicesGet)(pi_platform platform, pi_device_type device_type,
                             pi_uint32 num_entries, pi_device *devices,
@@ -305,6 +323,27 @@ pi_result OCL(piQueueCreate)(pi_context context, pi_device device,
   return cast<pi_result>(ret_err);
 }
 
+pi_result OCL(piextProgramConvert)(
+    pi_context context,  ///< [in] the PI context of the program
+    pi_program *program, ///< [in,out] the pointer to PI program
+    void **handle)       ///< [in,out] the pointer to the raw program handle
+{
+  // The PI program is the same as OpenCL program handle.
+  assert(program);
+  assert(handle);
+
+  if (*program == nullptr) {
+    // uninitialized *program.
+    assert(*handle);
+    *program = cast<pi_program>(*handle);
+  } else {
+    assert(*handle == nullptr);
+    *handle = *program;
+  }
+  cl_int result = clRetainProgram(cast<cl_program>(*handle));
+  return cast<pi_result>(result);
+}
+
 pi_result OCL(piProgramCreate)(pi_context context, const void *il,
                                size_t length, pi_program *res_program) {
 
@@ -992,6 +1031,7 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piPlatformsGet, OCL(piPlatformsGet))
   _PI_CL(piPlatformGetInfo, clGetPlatformInfo)
   // Device
+  _PI_CL(piextDeviceConvert, OCL(piextDeviceConvert))
   _PI_CL(piDevicesGet, OCL(piDevicesGet))
   _PI_CL(piDeviceGetInfo, clGetDeviceInfo)
   _PI_CL(piDevicePartition, clCreateSubDevices)
@@ -1019,6 +1059,7 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piMemRelease, clReleaseMemObject)
   _PI_CL(piMemBufferPartition, OCL(piMemBufferPartition))
   // Program
+  _PI_CL(piextProgramConvert, OCL(piextProgramConvert))
   _PI_CL(piProgramCreate, OCL(piProgramCreate))
   _PI_CL(piclProgramCreateWithSource, OCL(piclProgramCreateWithSource))
   _PI_CL(piclProgramCreateWithBinary, OCL(piclProgramCreateWithBinary))
diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp
index efce61d181a2d..b7b4a76c5a395 100644
--- a/sycl/source/detail/device_impl.cpp
+++ b/sycl/source/detail/device_impl.cpp
@@ -19,15 +19,30 @@ device_impl::device_impl()
     : MIsHostDevice(true),
       MPlatform(std::make_shared<platform_impl>(platform_impl())) {}
 
+device_impl::device_impl(device_interop_handle_t InteropDeviceHandle,
+                         const plugin &Plugin)
+    : device_impl(InteropDeviceHandle, nullptr, nullptr, Plugin) {}
+
 device_impl::device_impl(RT::PiDevice Device, PlatformImplPtr Platform)
-    : device_impl(Device, Platform, Platform->getPlugin()) {}
+    : device_impl(nullptr, Device, Platform, Platform->getPlugin()) {}
 
 device_impl::device_impl(RT::PiDevice Device, const plugin &Plugin)
-    : device_impl(Device, nullptr, Plugin) {}
+    : device_impl(nullptr, Device, nullptr, Plugin) {}
 
-device_impl::device_impl(RT::PiDevice Device, PlatformImplPtr Platform,
+device_impl::device_impl(device_interop_handle_t InteropDeviceHandle,
+                         RT::PiDevice Device, PlatformImplPtr Platform,
                          const plugin &Plugin)
     : MDevice(Device), MIsHostDevice(false) {
+
+  bool InteroperabilityConstructor = false;
+  if (Device == nullptr) {
+    assert(InteropDeviceHandle != nullptr);
+    // Get PI device from the raw device handle.
+    Plugin.call<PiApiKind::piextDeviceConvert>(&MDevice,
+                                               (void **)&InteropDeviceHandle);
+    InteroperabilityConstructor = true;
+  }
+
   // TODO catch an exception and put it to list of asynchronous exceptions
   Plugin.call<PiApiKind::piDeviceGetInfo>(
       MDevice, PI_DEVICE_INFO_TYPE, sizeof(RT::PiDeviceType), &MType, nullptr);
@@ -38,8 +53,10 @@ device_impl::device_impl(RT::PiDevice Device, PlatformImplPtr Platform,
       MDevice, PI_DEVICE_INFO_PARENT_DEVICE, sizeof(RT::PiDevice), &parent, nullptr);
 
   MIsRootDevice = (nullptr == parent);
-  if (!MIsRootDevice) {
+  if (!MIsRootDevice && !InteroperabilityConstructor) {
     // TODO catch an exception and put it to list of asynchronous exceptions
+    // Interoperability Constructor already calls DeviceRetain in
+    // piextDeviceConvert.
     Plugin.call<PiApiKind::piDeviceRetain>(MDevice);
   }
 
@@ -47,7 +64,7 @@ device_impl::device_impl(RT::PiDevice Device, PlatformImplPtr Platform,
   if (!Platform) {
     RT::PiPlatform plt = nullptr; // TODO catch an exception and put it to list
                                   // of asynchronous exceptions
-    Plugin.call<PiApiKind::piDeviceGetInfo>(Device, PI_DEVICE_INFO_PLATFORM,
+    Plugin.call<PiApiKind::piDeviceGetInfo>(MDevice, PI_DEVICE_INFO_PLATFORM,
                                             sizeof(plt), &plt, nullptr);
     Platform = std::make_shared<platform_impl>(plt, Plugin);
   }
@@ -75,13 +92,15 @@ cl_device_id device_impl::get() const {
     throw invalid_object_error("This instance of device is a host instance",
                                PI_INVALID_DEVICE);
 
+  const detail::plugin &Plugin = getPlugin();
   if (!MIsRootDevice) {
     // TODO catch an exception and put it to list of asynchronous exceptions
-    const detail::plugin &Plugin = getPlugin();
     Plugin.call<PiApiKind::piDeviceRetain>(MDevice);
   }
-  // TODO: check that device is an OpenCL interop one
-  return pi::cast<cl_device_id>(MDevice);
+  void *handle = nullptr;
+  Plugin.call<PiApiKind::piextDeviceConvert>(
+      const_cast<RT::PiDevice *>(&MDevice), &handle);
+  return pi::cast<cl_device_id>(handle);
 }
 
 platform device_impl::get_platform() const {
diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp
index 26bc59587b31a..71572216478d5 100644
--- a/sycl/source/detail/device_impl.hpp
+++ b/sycl/source/detail/device_impl.hpp
@@ -27,12 +27,20 @@ namespace detail {
 class platform_impl;
 using PlatformImplPtr = std::shared_ptr<platform_impl>;
 
+// TODO: SYCL BE generalization will change this to something better.
+// For now this saves us from unwanted implicit casts.
+struct _device_interop_handle_t;
+using device_interop_handle_t = _device_interop_handle_t *;
+
 // TODO: Make code thread-safe
 class device_impl {
 public:
   /// Constructs a SYCL device instance as a host device.
   device_impl();
 
+  /// Constructs a SYCL device instance using the provided raw device handle.
+  explicit device_impl(device_interop_handle_t, const plugin &Plugin);
+
   /// Constructs a SYCL device instance using the provided
   /// PI device instance.
   explicit device_impl(RT::PiDevice Device, PlatformImplPtr Platform);
@@ -196,7 +204,8 @@ class device_impl {
   is_affinity_supported(info::partition_affinity_domain AffinityDomain) const;
 
 private:
-  explicit device_impl(RT::PiDevice Device, PlatformImplPtr Platform,
+  explicit device_impl(device_interop_handle_t InteropDevice,
+                       RT::PiDevice Device, PlatformImplPtr Platform,
                        const plugin &Plugin);
   RT::PiDevice MDevice = 0;
   RT::PiDeviceType MType;
diff --git a/sycl/source/detail/program_impl.cpp b/sycl/source/detail/program_impl.cpp
index 6cd2f6b0380e8..a59f2414e11d0 100644
--- a/sycl/source/detail/program_impl.cpp
+++ b/sycl/source/detail/program_impl.cpp
@@ -80,17 +80,32 @@ program_impl::program_impl(
   }
 }
 
-program_impl::program_impl(ContextImplPtr Context, RT::PiProgram Program)
+program_impl::program_impl(ContextImplPtr Context,
+                           program_interop_handle_t InteropProgram)
+    : program_impl(Context, InteropProgram, nullptr) {}
+
+program_impl::program_impl(ContextImplPtr Context,
+                           program_interop_handle_t InteropProgram,
+                           RT::PiProgram Program)
     : MProgram(Program), MContext(Context), MLinkable(true) {
 
+  const detail::plugin &Plugin = getPlugin();
+  if (MProgram == nullptr) {
+    assert(InteropProgram != nullptr &&
+           "No InteropProgram/PiProgram defined with piextProgramConvert");
+    // Translate the raw program handle into PI program.
+    Plugin.call<PiApiKind::piextProgramConvert>(
+        Context->getHandleRef(), &MProgram, (void **)&InteropProgram);
+  } else
+    Plugin.call<PiApiKind::piProgramRetain>(Program);
+
   // TODO handle the case when cl_program build is in progress
   pi_uint32 NumDevices;
-  const detail::plugin &Plugin = getPlugin();
-  Plugin.call<PiApiKind::piProgramGetInfo>(Program, PI_PROGRAM_INFO_NUM_DEVICES,
-                                           sizeof(pi_uint32), &NumDevices,
-                                           nullptr);
+  Plugin.call<PiApiKind::piProgramGetInfo>(
+      MProgram, PI_PROGRAM_INFO_NUM_DEVICES, sizeof(pi_uint32), &NumDevices,
+      nullptr);
   vector_class<RT::PiDevice> PiDevices(NumDevices);
-  Plugin.call<PiApiKind::piProgramGetInfo>(Program, PI_PROGRAM_INFO_DEVICES,
+  Plugin.call<PiApiKind::piProgramGetInfo>(MProgram, PI_PROGRAM_INFO_DEVICES,
                                            sizeof(RT::PiDevice) * NumDevices,
                                            PiDevices.data(), nullptr);
   vector_class<device> SyclContextDevices =
@@ -109,16 +124,17 @@ program_impl::program_impl(ContextImplPtr Context, RT::PiProgram Program)
   SyclContextDevices.erase(NewEnd, SyclContextDevices.end());
   MDevices = SyclContextDevices;
   RT::PiDevice Device = getSyclObjImpl(MDevices[0])->getHandleRef();
+  assert(!MDevices.empty() && "No device found for this program");
   // TODO check build for each device instead
   cl_program_binary_type BinaryType;
   Plugin.call<PiApiKind::piProgramGetBuildInfo>(
-      Program, Device, CL_PROGRAM_BINARY_TYPE, sizeof(cl_program_binary_type),
+      MProgram, Device, CL_PROGRAM_BINARY_TYPE, sizeof(cl_program_binary_type),
       &BinaryType, nullptr);
   size_t Size = 0;
   Plugin.call<PiApiKind::piProgramGetBuildInfo>(
-      Program, Device, CL_PROGRAM_BUILD_OPTIONS, 0, nullptr, &Size);
+      MProgram, Device, CL_PROGRAM_BUILD_OPTIONS, 0, nullptr, &Size);
   std::vector<char> OptionsVector(Size);
-  Plugin.call<PiApiKind::piProgramGetBuildInfo>(Program, Device,
+  Plugin.call<PiApiKind::piProgramGetBuildInfo>(MProgram, Device,
                                                 CL_PROGRAM_BUILD_OPTIONS, Size,
                                                 OptionsVector.data(), nullptr);
   string_class Options(OptionsVector.begin(), OptionsVector.end());
@@ -137,12 +153,11 @@ program_impl::program_impl(ContextImplPtr Context, RT::PiProgram Program)
     MLinkOptions = "";
     MBuildOptions = Options;
   }
-  Plugin.call<PiApiKind::piProgramRetain>(Program);
 }
 
 program_impl::program_impl(ContextImplPtr Context, RT::PiKernel Kernel)
-    : program_impl(Context,
-                   ProgramManager::getInstance().getClProgramFromClKernel(
+    : program_impl(Context, nullptr,
+                   ProgramManager::getInstance().getPiProgramFromPiKernel(
                        Kernel, Context)) {}
 
 program_impl::~program_impl() {
diff --git a/sycl/source/detail/program_impl.hpp b/sycl/source/detail/program_impl.hpp
index f185b471df0a2..f8060a369de64 100644
--- a/sycl/source/detail/program_impl.hpp
+++ b/sycl/source/detail/program_impl.hpp
@@ -31,6 +31,11 @@ namespace detail {
 
 using ContextImplPtr = std::shared_ptr<detail::context_impl>;
 
+// TODO: SYCL BE generalization will change this to something better.
+// For now this saves us from unwanted implicit casts.
+struct _program_interop_handle_t;
+using program_interop_handle_t = _program_interop_handle_t *;
+
 class program_impl {
 public:
   program_impl() = delete;
@@ -71,21 +76,20 @@ class program_impl {
   program_impl(vector_class<shared_ptr_class<program_impl>> ProgramList,
                string_class LinkOptions = "");
 
-  /// Constructs a program instance from plugin interface interoperability
-  /// handle.
+  /// Constructs a program instance from an interop raw BE program handle.
+  /// TODO: BE generalization will change that to something better.
   ///
   /// The state of the constructed program can be either
   /// program_state::compiled or program_state::linked, depending on the state
-  /// of the ClProgram. Otherwise an invalid_object_error SYCL exception is
+  /// of the InteropProgram. Otherwise an invalid_object_error SYCL exception is
   /// thrown.
   ///
-  /// The instance of plugin interface program will be retained on
-  /// construction.
+  /// The instance of the program will be retained on construction.
   ///
   /// \param Context is a pointer to SYCL context impl.
-  /// \param Program is an instance of plugin interface interoperability
+  /// \param InteropProgram is an instance of plugin interface interoperability
   /// program.
-  program_impl(ContextImplPtr Context, RT::PiProgram Program);
+  program_impl(ContextImplPtr Context, program_interop_handle_t InteropProgram);
 
   /// Constructs a program instance from plugin interface interoperability
   /// kernel.
@@ -290,6 +294,9 @@ class program_impl {
   program_state get_state() const { return MState; }
 
 private:
+  // Deligating Constructor used in Implementation.
+  program_impl(ContextImplPtr Context, program_interop_handle_t InteropProgram,
+               RT::PiProgram Program);
   /// Checks feature support for specific devices.
   ///
   /// If there's at least one device that does not support this feature,
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 0d95d886b5061..d9719547c7a0c 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -433,7 +433,7 @@ RT::PiKernel ProgramManager::getOrCreateKernel(OSModuleHandle M,
 }
 
 RT::PiProgram
-ProgramManager::getClProgramFromClKernel(RT::PiKernel Kernel,
+ProgramManager::getPiProgramFromPiKernel(RT::PiKernel Kernel,
                                          const ContextImplPtr Context) {
   RT::PiProgram Program;
   const detail::plugin &Plugin = Context->getPlugin();
diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp
index 826ab66d8e187..e5ab759e76cae 100644
--- a/sycl/source/detail/program_manager/program_manager.hpp
+++ b/sycl/source/detail/program_manager/program_manager.hpp
@@ -66,7 +66,7 @@ class ProgramManager {
                                   const string_class &KernelName);
   RT::PiKernel getOrCreateKernel(OSModuleHandle M, const context &Context,
                                  const string_class &KernelName);
-  RT::PiProgram getClProgramFromClKernel(RT::PiKernel Kernel,
+  RT::PiProgram getPiProgramFromPiKernel(RT::PiKernel Kernel,
                                          const ContextImplPtr Context);
 
   void addImages(pi_device_binaries DeviceImages);
diff --git a/sycl/source/device.cpp b/sycl/source/device.cpp
index fa3ccddc55f02..9db8763602cf9 100644
--- a/sycl/source/device.cpp
+++ b/sycl/source/device.cpp
@@ -29,7 +29,7 @@ device::device() : impl(std::make_shared<detail::device_impl>()) {}
 
 device::device(cl_device_id deviceId)
     : impl(std::make_shared<detail::device_impl>(
-          detail::pi::cast<detail::RT::PiDevice>(deviceId),
+          detail::pi::cast<detail::device_interop_handle_t>(deviceId),
           *RT::GlobalPlugin)) {}
 
 device::device(const device_selector &deviceSelector) {
diff --git a/sycl/source/program.cpp b/sycl/source/program.cpp
index bb3011d614f3f..9f7c92e2bfaa3 100644
--- a/sycl/source/program.cpp
+++ b/sycl/source/program.cpp
@@ -30,7 +30,7 @@ program::program(vector_class<program> programList, string_class linkOptions) {
 program::program(const context &context, cl_program clProgram)
     : impl(std::make_shared<detail::program_impl>(
           detail::getSyclObjImpl(context),
-          detail::pi::cast<detail::RT::PiProgram>(clProgram))) {}
+          detail::pi::cast<detail::program_interop_handle_t>(clProgram))) {}
 program::program(std::shared_ptr<detail::program_impl> impl) : impl(impl) {}
 
 cl_program program::get() const { return impl->get(); }